simx: model the issue lock, per-slot dispatch queues and port-0 flush register as the RTL builds them - #423
Merged
Merged
Conversation
…h and DXA drain with the RTL The perf_gate and model_parity jobs timed out (run 36547574183). Since #416 the ALU/LSU/FPU block counts follow ISSUE_WIDTH, so every ISSUE_WIDTH=4 gate case elaborates a four-block model several times the size, and a hosted runner cannot build it within the job. Those cases now pin one block per unit; their cycles return to the pre-#416 values (wgmma-dxa-mcast nt4/nt16 reproduce them exactly). Two build costs amplified it: - rtlsim passed -O2 through CXXFLAGS, which lands after Verilator's per-class knobs and also optimizes the slow-path files. The level now goes through OPT_FAST/OPT_GLOBAL: wgmma-dxa 598 -> 224 CPU-min, same cycles. - blackbox built the driver with the caller's CONFIGS, then the run rebuilt it with the app's resolved CONFIGS. A runtime-<driver> target in each suite's common.mk builds it once, with the CONFIGS the run uses. SimX timing, both mismatches found on the DXA parity cases: - The dispatcher stepped through the issue-slot batches in fixed rotation, spending a cycle on every empty slot. VX_lane_dispatch grants the lowest batch holding an instruction and rests on the last when none does. - The DXA scatter drain (K-major, Flat, BlockMajor) gathered a word of elements per beat. VX_dxa_smem_wr writes one element per beat; DXA lmem_writes on wgmma-dxa now equal the RTL's. SimX against recorded RTL cycles, all 31 model_parity cases (RTL unchanged): wgmma-dxa 24.9% -> 5.5%, wgmma-sp-dxa 17.8% -> 5.2%, wgmma-dxa-mcast 9.8% -> 0.9%, sgemm2-fp16 4.5% -> 0.1%, wgmma-fedp2k-rs 4.0% -> 0.7%; every other case unchanged. The multicast case drops its known_issue and 10% tolerance, the sparse case its 10% tolerance; the two remaining known_issue reasons now state the measured gap. The wgmma-dxa residual is C-tile store drain at CTA turnover. Not run here: the full rtlsim model_parity and perf_gate gates. perf baselines are stale (fingerprint) and need --update-baselines. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…the RTL does SGEMM on two cores (model_parity sgemm-mc) ran 10.5% slower in SimX than in the RTL. The kernel's loop is a dependent eight-FMA chain with thirteen loads, so it is bound by the FPU's two tags. Three SimX timing paths diverged from the RTL: - Issue: VX_scoreboard drops a warp from arbitration while its target FU's dispatch queue is going-full. SimX only de-prioritized such warps and let them all through when every ready warp targeted a congested FU. Excess FPU ops (7 issued and not yet in the FPU, against the RTL's 4) filled the two shared operand collectors and blocked loads and ALU ops behind them. The warps fell into lockstep and their load writebacks committed in bursts of 17-40 beats that starved the lower-priority FPU. The gate is now hard. The suppress-mask arbiter overloads lose their last caller and are removed. - FPU: VX_fpu_unit holds each tag until the result fires into the commit arbiter (no buffer between them at full bandwidth). SimX freed tags on a timer, buffered extra results and added two cycles to every latency. The response port now has one slot per tag, so taking a result off it frees the tag. At full bandwidth the commit arbiter reads the port directly. With partial bandwidth the lane-gather stage does. The latencies are the configured unit latencies. - Crossbar: TxCrossBar ignored its arbiter type and always granted the lowest input. Round-robin outputs now rotate from the last accepted input, reset to the last input, as VX_stream_xbar / VX_rr_arbiter do. The cache crossbars (round-robin) change; local memory (priority) does not. SimX against the recorded RTL cycles, all 31 model_parity cases: seven were over 5%, three are now. sgemm-mc 10.49% -> 2.06%, sgemmx 3.93 -> 0.53, rt_raycast 6.27 -> 2.83, fp16 6.28 -> 3.32, fp16-mc 4.07 -> 1.47, wgmma-sp-dxa 5.15 -> 4.21. Three got worse, all from the issue gate: sgemv 7.23 -> 9.59, wgmma-dxa 5.46 -> 8.53, mxfp8 0.82 -> 3.36. The gate matches the RTL, so these expose other SimX timing errors it was offsetting. They are not localized yet. Catalog: sgemmx and fp16 return to the 5% default and wgmma-sp-dxa loses its known_issue. sgemv and copy-mcast keep their existing 10% tolerance, and the wgmma-dxa reason states its measured gap. Verified: SimX-only sweep of all 31 model_parity cases against recorded RTL cycles (the RTL is unchanged). The final channel-based FPU wiring reproduces the swept cycles exactly on sgemm-mc, sgemv, fp16 and wgmma-dxa. sgemmx moves 2 cycles because the event wheel's push_back does not preserve FIFO order (fixed separately). FP-heavy apps pass in both FPU gather modes. Not run: the full pytest model_parity suite with its rtlsim legs. perf_gate runs on rtlsim only and cannot move. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The RTL memory scheduler takes a response-buffer entry when the core request is accepted into its queue and frees it at the last response, so loads waiting in the queue and loads in flight share one pool of LSU_PENDING_SIZE. SimX allocated the pending entry only when the request left req_queue, so up to four more reads could sit in the LSU: traces of parity-sgemv show 12 loads inside the SimX LSU against the RTL's 8, and each load waited correspondingly longer (matrix-row loads 59 cycles from LSU entry vs RTL 50). Reads (loads and AMOs) now count against the pool on entry and turn the reservation into their tag on issue; stores and fences are unaffected. Load latency from LSU entry now tracks the RTL (48.4 vs 50.0 and 34.1 vs 31.5 for the two sgemv streams), and parity-sgemv drops from 9.6% to 0.65%, so its 10% tolerance goes back to the 5% default. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…e escapes Remove the last known_issue marker (model_parity-wgmma-dxa) and every per-case parity tolerance override (copy-mcast at 10%, the redundant 5% on sgemm-mc), so each model_parity case asserts cycle agreement within the 5% default and a case over the threshold fails CI instead of being reported as an expected failure. The catalog now carries no known_issue anywhere. At the committed SimX model this turns two cases red: wgmma-dxa (SimX 8.5% slower than the RTL) and copy-mcast (6.8%). The pending SimX tile-buffer and LMEM DMA port model brings wgmma-dxa to 2.4%; copy-mcast needs the SimX per-stage latency rework. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…lds them wgmma-dxa ran 8.5% slower in SimX than the RTL. The four warps of a warp-group interleaved their C-tile stores one instruction at a time, so the stores evicted each other's lines and data-cache bank 0 fetched every C line two or three times. In the RTL the warps leave the MMA staggered and store back to back. The stagger comes from the TCU's shared B buffer. It holds one bank row, keyed by the lowest block with a WGMMA compute uop, and a block whose (desc_b, step_k, step_n) differs waits; each A buffer holds one k-stripe. SimX planned every A and B line of the WGMMA up front and held all blocks until all were resident, so the blocks ran in lockstep. The tile buffers now refill per key with the RTL's request counts, one read in flight per buffer and fixed-priority arbitration, and a setup uop drops them at the end of its cycle. That alone moved wgmma-dxa only to 8.2%: a refill took 4 cycles in SimX against 1 in the RTL. SimX routed TCU reads and DXA writes through the LSU's per-word bank crossbar, behind the LSU ports and with the TCU ahead of the DXA. Local memory now has the RTL's row-wide DMA port: DXA first, then TCU, one access per cycle, priority over the LSU on the banks it touches, and read data the cycle after the request. SimX vs RTL cycles (model_parity, all 31 cases swept against recorded RTL cycles): wgmma-dxa 8.53% -> 2.43%, wgmma-dxa-mcast 2.26% -> 0.20%, wgmma-sp-dxa 4.21% -> 0.08%, wgmma-int8 0.15% -> 0.32%, wgmma-fedp2k-rs 0.19% -> 0.30%; every other case bit-identical. All 47 functional dxa and tensor_wg SimX cases pass. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
copy-mcast ran 6.8% slower in SimX. Measured per opcode, every gap was in the ops that hold their warp until they resolve: a branch took 18 cycles from schedule to the warp's next schedule against the RTL's 17, SPLIT and WSYNC 20 against 17, JOIN 20 against 18. The RTL releases such a warp from the unit that resolves it, through one register (a join through one more, for the divergence-stack pop); SimX released it only when the result left the functional unit. The ALU and SFU now send the release to the core on a one-cycle channel, and a TMC that disables every thread retires its warp the same way. Releasing on time exposed three places where SimX ran faster than the RTL, compared per event on vecadd: - a warp freed by TMC was reused for the next CTA in the same cycle; the RTL picks from a registered active set, registers the pick, and the new warp is schedulable the cycle after; - integer multiply/divide had the plain ALU's latency; the multiplier pipeline is two stages longer; - a dependent instruction dispatched 4 cycles after its producer committed, against 7 in the RTL, while ALU and SFU results reached commit 2 cycles late. The scoreboard now clears a destination two cycles after commit, and the ALU and SFU result latencies lose the matching two cycles. SimX vs RTL cycles, all 31 model_parity cases (on top of the previous commit): copy-mcast 6.83% -> 2.22%, softmax 3.44% -> 0.04%, copy 4.06% -> 0.35%, sgemm2-dxa 2.83% -> 0.81%, vecadd 2.94% -> 1.54%; 30 of 31 within 5%. Known regression: parity-sgemm-mc goes from 2.24% to 11.85% (SimX slower) and fails the 5% gate. Most of it follows the scoreboard release delay (5.11% without it). On a reduced run FMADD dispatch-to-commit averages 49 cycles in SimX against 15 in the RTL: SimX dispatches the FMA chain early and it queues in front of the FPU and at commit, where the RTL holds it in the instruction buffer. Under investigation; raycast (4.53%), draw3d (4.73%) and sparse-fp16 (4.89%) are close to the limit. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
… register as the RTL builds them Brings every model_parity case within the 5% default. The three cases over it: sgemv 6.38% -> 0.83%, wgmma-fedp2k-rs 6.21% -> 0.92%, wgmma-dxa-mcast -7.31% -> 3.32%. Each change models an RTL structure confirmed in the RTL source and against debug traces from both models. Issue and dispatch: - Issue lock: a warp inside a locked uop sequence (WGMMA, RTU TRACE, OM export) blocks every other warp of its issue slot, whatever the unit. The lock reopens when the last uop issues. SimX used to lock only the target unit and released the lock later, when the unit accepted the last uop. - Dispatch queues: one per issue slot, with the lowest slot served first. A slot's issue credit returns when its op leaves the queue (new ReleaseOut channel). Dispatchers bind straight to the unit inputs. SimX used to share one queue per unit and return the credit when the unit accepted the op. The queue order sets how CTAs split across warps in wgmma-fedp2k-rs. - The near-full flag reaches issue through two registers, and a refused operand-to-dispatch hand-off counts as that unit's stall. - The uop sequencer offers a macro's first uop one cycle after the macro reaches the buffer front, timed from each warp's issue history rather than instruction uuids. Release builds zero the uuids, so a uuid-based check timed differently in release and debug. TCU and pack-load macros no longer stall fetch. Units and memory: - Reduced-width ALU/FPU: one extra input register and one extra lane-gather output register. - Multiply/divide latency is 3; one result per cycle is kept. - TCU admission holds while a completing result finds an earlier one still waiting for commit. - The first data-cache port also carries cache-flush injection, so its requests take one more register. Same-bank pairs are therefore served port 1 first, which sets the eviction order under the cache thrash in sgemv. Commit and scoreboard: - The fflags/frm CSR fields are scoreboarded like registers. - A result crosses one more register before the commit arbiter. - Scoreboard release delay goes from 2 cycles to 1. FPU latency coupling (VX_config.toml): - The DPI FPU is a fast simulation stand-in for STD, so it now takes STD's FDIV/FSQRT latencies from the same toml helpers (17/17, or 32/32 with D) instead of its own 15/10. Changing STD's latency therefore changes DPI's too. DSP and FPNEW keep their own values. - SimX (STD) and rtlsim (DPI) now resolve to the same latencies. Perf baselines, regenerated with `pytest -m perf_gate --update-baselines`: - The latency change moves only kernels that divide or take square roots, and by little: raycast +0.17%, stencil3d +0.18%, draw3d +0.09%, softmax -0.01%, rt_raycast -0.14%. - The other moves were already present before this change: upstream RTL commits merged after the last recording, and catalog configs changed without a re-record. wgmma-dxa (-13.9%) and wgmma-dxa-mcast (+3.8%) were outside the 2% tolerance, so perf_gate was already failing on them. Verified on a 32-bit build: - model_parity: 31/31 within 5%, worst fp16 -4.51%. The fp16 margin is only partly explained. - perf_gate: 49/49 against the new baselines. - All 23 SimX CI test categories pass: 261/261. - Release and debug SimX builds give identical cycle counts on wgmma and draw3d. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…put change GitHub cannot condition a cron schedule on the repository. Both gates therefore started a run every night just to find out whether master had moved. On a quiet night that left an empty run behind. And because the check keyed on master's SHA, a docs-only or SimX-only commit triggered a full multi-hour sweep. ASIC gate (hosted runners, no host to keep state on) now arms and disarms itself: - The workflow stays disabled between qualifying pushes, so its schedule does not fire. - ci.yml's plan job enables it on a push to master that touches hw/, the configuration TOMLs, VERSION, the gate's scripts, catalog or baselines, or the workflow itself. An indeterminate diff also arms it. - The scheduled run disables the workflow as its first step, so a push that lands mid-sweep re-arms the next night. A build error re-enables it to retry. A pass or a regression is a verdict, and its run is the record. - The SHA-marker cache and the skipped-night republish are gone. FPGA gate (self-hosted Vivado runner) loses its schedule: - A nightly cron on the runner's host runs ci/fpga_gate_dispatch.sh. It dispatches the workflow only when a synthesis input changed since the last gated SHA (the workflow's existing state file), the runner is up, and no gate run is queued or in progress. - The workflow drops its skip check, its republish step and the force input. A dispatched run always has work. synth_report.py loses --note: only the removed republish steps used it. Host setup still needed: install the crontab line from the script's header, and authenticate gh with a token that can dispatch workflows. Not run: these workflows only execute on GitHub. The dispatcher passed bash -n but has not been exercised end to end. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Brings every
model_paritycase within the 5% default tolerance. No tolerance is widened and no case is markedknown_issue.Every change makes SimX copy an RTL structure that was checked against the RTL source and against debug traces from both models.
Positive means SimX is slower than the RTL.
This branch carries 7 commits; the last one is the subject of this PR. The 6 before it are earlier parity work from local master that was never pushed, and the last commit builds on them.
Changes
Issue lock and dispatch queues
Memory
Smaller timing fixes
FPU latency coupling (
VX_config.toml)Perf baselines
pytest ci -m perf_gate --update-baselines.wgmma-dxa(−13.9%) andwgmma-dxa-mcast(+3.8%) were already outside the 2% tolerance.Test plan (32-bit)
pytest ci -m model_parity: 31/31. Worst case fp16 at −4.51%; that margin is only partly explained.pytest ci -m perf_gate: 49/49 against the regenerated baselines.🤖 Generated with Claude Code