Skip to content

CHANGELOG.md

Latest commit

 

History

History
111 lines (99 loc) · 40.2 KB

CHANGELOG.md

File metadata and controls

111 lines (99 loc) · 40.2 KB
Raw

Changelog

All notable changes to Vortex are documented here. The format is based on Keep a Changelog, and the project follows the version pins recorded in VERSION (VORTEX_VERSION, TOOLCHAIN_REV, GEM5_REV).

[Unreleased]

Fixed

  • Verilator rolled back to 5.050 — 5.052 cannot build the rtlsim model. On the smallest configuration in the suite (config1:diverge-1, 1 warp / 1 thread) verilator_bin itself is OOM-killed (%Error: Verilator threw signal 9, Error 137). Peak RSS for that build on one machine with identical flags and -j: 5.046 → 0.51 GB, 5.050 → 0.54 GB, 5.052 → 204.25 GB — a 400x blowup in verilator's own elaboration, not in the C++ compile. A 16 GB hosted runner never survives it, so every rtlsim job died with exit 143 as soon as CI reached the test stage. 5.050 is the newest release that works, isolating the regression to the single 5.050 → 5.052 step. Consumer-side coupling: 5.050 packages as three parts where 5.052 was two, so ci/toolchain_install.sh.in specifies {a..c} for ubuntu/focal; the published parts and that count must move together or --verilator 404s.
  • Yosys 0.69 deadlocked on every ABC pass. 0.69 gates a persistent ABC process pool behind YOSYS_ENABLE_READLINE: it posix_spawns yosys-abc -s, writes empty\nsource <script> to its stdin and reads until a YOSYS_ABC_DONE sentinel. ABC runs the script and writes output.blif, but yosys never returns from its read — yosys blocks in pipe_read, ABC in select() on a stdin yosys will not close — so every Yosys-flow build sat until its timeout with no output. Reproduces on a three-line Verilog module. The prebuilt is now built with -DYOSYS_WITHOUT_READLINE=ON -DYOSYS_WITHOUT_EDITLINE=ON, which leaves REUSE_YOSYS_ABC_PROCESSES undefined and restores the classic abc -s -f path; readline only provides line editing in the interactive shell. New ci/yosys_install.sh.in records the build and asserts it with a smoke test that exercises the ABC pass — validated against the known-bad build, which it rejects.
  • Verilator 5.052 wide-signal ports broke the opaesim build. 5.052 emits wide AFU ports as VlWide<> objects, which no longer decay to void*, so four memcpy sites in sim/opaesim/opae_sim.cpp failed to compile. They now go through the VDataCast helper in sw/common/util.h that sim/rtlsim/processor.cpp already used; the helper resolves for both layouts, so the fix builds against 5.046 and 5.052 alike.
  • The Yosys flow wrote an unclean netlist. hw/syn/yosys/run_synth.sh never ran opt_clean after mapping, so the netlist carried ~129 undriven assign w = 'hx wires left over from sv2v function scopes, and stat counted dead cells as area. Harmless under Yosys 0.57, fatal under 0.69: it marks those wires signed, which OpenSTA's Verilog reader rejects outright (Error: 171 ... syntax error), so timing and power produced nothing at all and the DUT reported no Fmax. Why it looked like a tool incompatibility: the debris was in both toolchains' netlists all along — only the signed qualifier was new.
  • A single signed net made OpenSTA reject the whole netlist. OpenSTA's Verilog grammar carries no signed keyword, so one signed net declaration fails the read outright (Error: 171 <netlist> line N, syntax error) and the DUT reports no timing or power — the rtu DUT hit this on a live FMA net. Yosys preserves net signedness as of 0.69 where 0.57 dropped it, so the netlist, not the RTL, changed. hw/syn/yosys/run_synth.sh now strips the qualifier from declarations in the mapped netlist before STA reads it; signedness is information-free between library cells, so nothing STA computes can move. This is a workaround, not a fix (AGENTS.md §3) — the defect is in OpenSTA's reader, which should accept and ignore signed, and the follow-up is a patch there.
  • Flop mapping cost 25% of sequential area under Yosys 0.69. dfflibmap picks one library cell per FF type: 0.57 took the cheapest match, 0.69 prefers a non-inverting output. ASAP7 carries a non-inverting flop only at x4 drive (DFFHQx4, 0.3645) against DFFHQNx1 (0.2916), so all 129,564 flops in the tex DUT moved to x4 — +9% total cell area, +25% sequential area, with Fmax slightly worse and reported power collapsing 73% as vectorless activity propagation changed with the flop output polarity. New DFF_DONT_USE knob (run_synth.sh, defaulted per PDK in hw/syn/yosys/Makefile) withholds those cells, which is the standard dont-use-list mechanism; tex returns to within ±1% of its 0.57 Fmax, area, cell count and power.

[3.0] — 2026-06-08

The 3.0 release introduces a fixed-function graphics stack (rasterizer, texture units, and output mergers), tensor core structured sparsity (2:4), warpgroup-level matrix multiplication (WGMMA), global-to-local data transfer acceleration (DXA), a new hardware kernel scheduler (KMU) and Command Processor (CP) architecture, a new asynchronous runtime API (vortex2.h), asynchronous barriers with arrive/wait/event semantics, compressed instruction set (RVC) support, hardware atomics, an MMU/SV32 virtual memory stack, a Mesa/lavapipe Vulkan backend (vortexpipe), HIP via chipStar, gem5 integration, a SimX v3 TLM architecture with fixed-size handshake channels, productized Synopsys and Yosys ASIC synthesis flows, and a refreshed toolchain (LLVM 20, POCL 7.0). Build and configuration infrastructure was reworked: TOML-driven HW configuration (VX_config.toml + VX_types.toml) decoupling SimX/runtime from the RTL source tree, a VX_CFG_ macro namespace that resolves toolchain preprocessor collisions, retirement of the global toolchain_env.sh to enable parallel multi-version Vortex worktrees on the same shell, consolidation of kernel//runtime/ under a shared sw/ root, a single-source VERSION file driving CI toolchain pinning, Perfetto trace export (ci/perfetto.py), and new top-level AGENTS.md + CONTRIBUTING.md for AI-agent and contributor workflows.

Added

  • FireSim FPGA-accelerated simulation (Alveo U55C). Vortex runs as a host-decoupled FireSim target — Chisel BlackBox + FASED memory timing model (hw/syn/firesim/src/main/scala/VortexTarget.scala), AFU wrapper (hw/rtl/afu/firesim/VX_firesim_wrap.sv), transport (sim/firesim/firesim_sim.cpp) and runtime HAL (sw/runtime/firesim/vortex.cpp) — selected by --driver=firesim, built from vortexgpgpu/firesim vortex_3.x consumed as $(FIRESIM_PATH) and packaged by ci/toolchain_{install,prebuilt}.sh --firesim; the image is $(FPGA_BIN_DIR)/firesim.xclbin, and XCLBIN_PATH now names the image for both the xrt and firesim drivers, replacing XRT_XCLBIN_PATH. Why: every target clock edge is scheduled by the host and target DRAM is served by a timing model, so cycle counts are reproducible run to run and the modelled memory system can be varied independently of the card — sgemm -n1024 (1.19 B instructions) runs in ~85 s with instruction counts matching rtlsim exactly, at 10.09 % LUT / 5.41 % BRAM / 0.27 % DSP and WNS +5.283 ns. Limits: U55C only — the host-memory base address, the HBM[0:7] binding and the single memory bank are fixed to that card, and on any other platform the target DRAM path fails silently. See docs/designs/firesim_integration.md.
  • AMD Alveo V80 support (aved backend). New AFU, runtime, Verilator model and synthesis flow targeting the V80 through the SLASH platform (VRT runtime, slashkit linker, IP-XACT packaging) — the board is not an XDMA shell, so XRT does not apply. TARGET=hw|sim|avedsim selects what executes the RTL; avedsim is the iteration loop and matches the xrt backend cycle-for-cycle.
  • Host-resident Command Processor ring on V80. The CP pulls commands from coherent host memory over the QDMA slave bridge, as mainstream GPU command processors do, backed by a new host-buffer allocator spanning the SLASH kernel module, libslash, vrtd and VRT.
  • TCU tfr arithmetic backend. New in-house, fully-synthesizable fused dot-product running integer and floating-point through one shared 4-cycle pipeline; gated by VX_CFG_TCU_TYPE_TFR. ** Adds FP8 (e4m3), BF8 (e5m2), and TF32 on top of the v2.x set (fp32 / fp16 / bf16 / i32 / i8 / u8 / i4 / u4). Each is gated by its own VX_CFG_TCU_{FP8,BF16,TF32}_ENABLE; format dispatch is unified across all four FEDP backends.
  • Tensor-core structured sparsity (2:4). VX_tcu_sp_mux + VX_tcu_sp_meta datapath plus host compress_2to4_matrix / prune_2to4_matrix helpers; gated by VX_CFG_TCU_SPARSE_ENABLE.
  • Warpgroup-level MMA (WGMMA). Per-warp NRA=4 / variable-NRC fragment layout, S/R source modes, smem descriptor path; gated by VX_CFG_TCU_WGMMA_ENABLE.
  • Data-transfer Acceleration (DXA). Async global→local DMA engine for tile staging (hw/rtl/dxa/ + sim/simx/dxa/).
  • Hardware Kernel Management Unit (KMU). New scheduler block (hw/rtl/VX_kmu.sv + sim/simx/kmu/) that owns CTA dispatch from the CP launch path.
  • Launch-level CTA clustering. vx_launch_info_t::cluster_dim[3] (sw/runtime/include/vortex2.h:202) guarantees every K-CTA group is co-resident on one core; KMU iterates intra-cluster offsets first (sim/simx/kmu/kmu.cpp:63) and the CTA dispatcher reserves K contiguous LMEM slots so DXA Path A multicast can target issuer + r*stride.
  • Command Processor (CP) v3. New hw/rtl/cp/ block + host-resident command ring (CMD_LAUNCH, CMD_MEM_*, CMD_DCR_*, CMD_CACHE_FLUSH, CMD_EVENT_*); integrated end-to-end across xrt, opae, simx, rtlsim behind VORTEX_USE_CP.
  • Asynchronous vortex2.h runtime API. Queues, events, modules, kernels, UVA raw-pointer kernel args, per-queue worker thread; legacy vortex.h retained as a thin wrapper.
  • C++ software CP model (sim/common/cmd_processor.cpp) shared by simx and rtlsim.
  • Graphics stack (RASTER / TEX / OM / RTU). Fixed-function 3D pipeline: hw/rtl/{raster,tex,om,rtu}/ + VX_graphics.sv + matching SimX models; --graphics regression group.
  • Public host-side graphics API. sw/runtime/include/graphics.h exposes vortex::graphics::Binning() (triangle setup + tile binning producing the on-wire rast_prim_t stream the RASTER unit reads) plus self-contained vertex_t / primitive_t input types and DCR address helpers for external Vulkan/HIP/OpenGL drivers.
  • Canonical on-wire graphics ABI in sw/kernel/include/vx_graphics.h. Templated POD vortex::graphics::fixed_t<F> (Q15.16 / Q?.24 with full arithmetic, all members public + trivially copyable) plus the on-wire structs (vec3e_t, rast_prim_t, rast_attribs_t, rast_tile_header_t, etc.) and 8888 pixel helpers.
  • Vortex SDK install layout. make install produces $VORTEX_PATH/{kernel,runtime}/{include,lib<XLEN>} plus lib/pkgconfig/vortex-{runtime,kernel}.pc (auto-generated from sw/runtime/vortex-runtime.pc.in and sw/runtime/vortex-kernel.pc.in at configure time). Default prefix is <build>/install. Why: gives downstream tools (mesa-vortex, pocl-vortex, chipStar) a single $VORTEX_PATH env var + pkg-config integration shape.
  • Vulkan support via a new Mesa Gallium driver vortexpipe selected through the lavapipe ICD; tests/vulkan/ suite (compute, draw3d, depth, textured, raytrace); Mesa shipped via the prebuilt toolchain; rv64 path enabled.
  • HIP support on rv32 + rv64 via chipStar. chipStar's hipcc now accepts --offload-pointer-width={32,64} and emits SPIR-V with the matching OpMemoryModel Physical{32,64}; a single libCHIP.so ships both widths' rtdevlib modules and selects at runtime via CL_DEVICE_ADDRESS_BITS. POCL on rv32 Vortex accepts the resulting Physical32 SPIR-V cleanly. See docs/designs/hip_on_vortex_chipstar.md.
  • OpenCL image support on PoCL-Vortex. clCreateImage, read_image{f,i,ui} / write_image{f,i,ui}, get_image_*, and samplers on the Vortex PoCL device (CL_DEVICE_IMAGE_SUPPORT). Sampling is classified per read into two tiers: a fixed-function TEX fast path (Tier A) that routes an FF-representable read_imagef (2D, power-of-two, 8-bit UNORM RGBA/BGRA, nearest/bilinear, clamp/repeat/mirror) through the existing vx_tex4 hardware sampler on a host-bound TEX stage — zero new silicon, ~46 % fewer instructions for nearest and ~4.7× for bilinear — and a permanent scalar software sampler (Tier C) for every other format/mode, chosen by a gated fallback (never a silent wrong result). The TEX-stage binds and the kernel launch are submitted as one CP command batch (vx_enqueue_commands, a single doorbell), the compute-path analog of CMD_DRAW. Image builtins are Vortex-specific scalar (the no-vector backend cannot select the generic vector image builtins). Tier equivalence (hardware vs software across both TEX stages, the software fallback, and every wrap/filter mode) is asserted by tests/opencl/image_hwsw; see also tests/opencl/image_{ff,ff_linear,ff_bgra,copy,linear,int,get}. Why: completes the OpenCL 1.2 feature set and moves the compute path onto the FF-reuse + CP-batch posture the graphics stack already holds. The image builtins + driver live in the PoCL-Vortex tree (vortex_3.x); the runtime ships in the prebuilt toolchain.
  • Fix: FF TEX clamp-to-edge for negative coordinates. hw/rtl/tex/VX_tex_sat.sv mis-saturated a negative texture coordinate to the high edge (its sign-extension bits falsely tripped the overflow path); the overflow term is now sign-gated, so negatives clamp to texel 0, matching the SimX TextureWrap model. Surfaced by OpenCL read_imagef with clamp + out-of-range coordinates (ordinary texture mapping stays in range, so the graphics TEX tests are unaffected).
  • Hardware atomics. RISC-V A-extension (LR/SC reservation table + cache-resident AMO* RMW), gated by VX_CFG_EXT_A_ENABLE. AMOs complete at the LLC while non-LLC banks invalidate the line on passthrough, so atomics are correct across the full L1/L2/L3 cache hierarchy.
  • In-house IEEE-754 FPU (VX_fpu_std), now F32 and F64. Fully RV-compliant scalar FPU built from Vortex-owned blocks, covering both single (F) and double (D) precision natively — VX_fma_unit (separate F32/F64 fused multiply-add cores), merged-format VX_fdivsqrt_unit (radix-2 non-restoring FDIV + FSQRT; one carry-save datapath sized for the widest format — 17-cycle F32-only, 32-cycle when D is enabled), merged-format VX_fcvt_unit (I2F/F2I/F2F incl. FCVT.S.D subnormal/overflow), and merged VX_fncp_unit (sign-inject/min-max/compare/class/move with F32 NaN-box checking) — selected via VX_CFG_FPU_TYPE_STD. Validated against the full rv64u[fd]-p riscv-tests ISA suite on the RTL FPU. Why: removes the FPNEW dependency entirely — FPNEW is no longer required for D support on any flow (ASIC/Yosys/Synopsys and FPGA), only optionally selectable; the native units deliver higher fmax, lower latency, and smaller area, and the in-tree source unblocks block-level tuning that vendoring made impractical.
  • RISC-V Zicond (conditional ops). CZERO.EQZ / CZERO.NEZ integrated end-to-end (decode in VX_decode.sv, ALU in VX_alu_int.sv); gated by VX_CFG_EXT_ZICOND_ENABLE. Adds an ISA-level branchless-select primitive used by LLVM 20's codegen.
  • Pack-load intrinsics (vx_packlb_f / vx_packlh_f). Single-instruction strided loads that fold 4×byte (PACKLB) or 2×halfword (PACKLH) loads into one front-end issue, expanded by VX_uop_packld into N back-to-back LSU uops with eff_rs1 = rs1 + rs2 × uop_idx. Used heavily by sw/kernel/include/vx_tensor.h for TCU tile-row packing.
  • Wallace-tree + folded-radix multipliers (VX_wallace_mul, VX_fold_mul). New hw/rtl/libs/ multiplier blocks. Why: used by VX_fma_unit (mantissa multiply) and the TFR TCU backend (per-format integer / fp multiply); shared structural multiplier lets both blocks pick the area/latency trade-off via a single point of change.
  • Kogge-Stone parallel-prefix adder (VX_ks_adder). Logarithmic-depth carry-propagate adder under hw/rtl/libs/. Why: drives the long carry chains in VX_fma_unit (exponent/mantissa alignment + final add) and the TFR TCU's FEDP final accumulator without the ripple-carry timing penalty Verilator's default + emits.
  • Stream split/join primitives (VX_stream_dispatch, VX_stream_fork, VX_stream_join). New hw/rtl/libs/ modules consumed by VX_dcr_arb, VX_dxa_dispatch, and VX_gbar_arb. Why: replaces ad-hoc 1→N and N→1 ready/valid plumbing duplicated across the DCR, DXA, and global-barrier paths with one tested, parameterizable primitive; cuts ~3 copies of the same handshake state machine.
  • Inference-based integrated clock gating (VX_clockgate). Synthesizable ICG cell at hw/rtl/libs/VX_clockgate.sv; instantiated for per-core gating in VX_socket.sv:407. Why: gives the synthesis tools a single recognizable ICG inference pattern so ASIC flows (Synopsys/Yosys) generate proper latch-based gates instead of glitching AND-based gating; per-core gating is the first power-domain leverage point.
  • Cache fill-forwarding (FILL_FORWARD). The read-prefix of a fill's MSHR pending chain completes directly from the staged fill sector instead of replaying through the bank pipeline (hw/rtl/cache/VX_cache_bank.sv, SimX mirror in sim/simx/mem/cache.cpp); the first write/AMO entry falls back to the replay path so coalesced same-line accesses keep program order. Why: standard GPU L1 practice (fill-buffer forwarding) — pending misses stop paying the replay round-trip and the input arbiter stays open to demand requests while a chain drains (sgemm −2.7%, sgemv −5.0%, sgemm-mc −2.2% cycles); also hardens the tag-store read-first bypass and MSHR chain linking against the request-after-fill races the old replay serialization masked.
  • Compressed instruction set (RVC). New VX_decompressor block in the fetch stage (hw/rtl/core/VX_decompressor.sv + sim/simx/decompressor.cpp); gated by VX_CFG_EXT_C_ENABLE. v2.x shipped the test binaries but had no decompressor.
  • Asynchronous barriers with arrive / wait / expect_tx semantics. VX_bar_unit + vortex::barrier host API; expect_tx is the hook DXA multicast uses to declare expected bytes.
  • MMU / virtual memory (SV32, rv32-only). Host-shadow page table + DeviceMemIO refactor + --vm regression group.
  • GEM5 integration. VortexGPGPU SimObject + x86/aarch64 host runtimes; ci/regression.sh --gem5 + VORTEX_GEM5_ARM=1.
  • SimX v3 TLM architecture. Transaction-level memory packets (MemReq/MemRsp with shared_ptr<mem_block_t> payloads) and reusable TLM cache / switch / coalescer modules across the L1/L2/L3 hierarchy, tcache/ocache/rcache, DXA, and CP DMA paths.
  • ASIC synthesis flows. hw/syn/{synopsys,yosys}/ productized: shared hw/syn/common.mk, bundled NanGate_15nm_OCL.db standard cells, standardized OPT_LEVEL; legacy hw/syn/modelsim flow retired.
  • Synopsys multi-PDK support. hw/syn/synopsys/Makefile now exposes three target PDKs via LIB_TGT selection: ASAP7 (7nm), SAED14 (14nm SLVT), and NanGate 15nm OCL (default). Per-PDK SRAM mappings, pin polarity, and address packing handled in hw/syn/synopsys/project.tcl (SRAM_PINS + family dispatch). Why: lets users compare PPA on a real foundry-style 7nm/14nm flow without rewriting the synthesis scripts per technology node.
  • Yosys + OpenSTA end-to-end ASIC pipeline. hw/syn/yosys/run_synth.sh runs synthesis → tech-map → area (stat -liberty) → SRAM-cost estimation (sram_cost.py) → OpenSTA timing + power (run_sta.tcl, gated by RUN_STA=1). Makefile targets: synth, techmap, timing. Why: gives the open-source flow the same area/timing/power deliverables as the Synopsys flow — no commercial-tool license needed for first-pass PPA exploration.
  • SAIF switching-activity workflow for power analysis. ci/blackbox.sh --saif [--saif_file=...] captures gate-level activity from rtlsim runs; both the Synopsys (project.tcl:489-922 — read_saif -auto_map_names + report_saif) and Xilinx (hw/syn/xilinx/dut/common.mk:31) power flows consume the resulting SAIF for vector-driven power estimation. A matching --vcd [--vcd_file=...] switch gates waveform output independently of --debug, so trace, waveform, and activity capture compose freely. Why: replaces vectorless power estimation with activity-annotated numbers tied to a real workload, dropping the order-of-magnitude error band that purely static analysis carries.
  • OpenSTA tool packaged with the toolchain. New sta() function in ci/toolchain_install.sh.in + ci/toolchain_prebuilt.sh.in; installed under $TOOLDIR/sta/ and consumed by both the Yosys flow (run_sta.tcl) and hw/syn/common.mk as $(STA). Why: makes OpenSTA a first-class peer of Verilator/sv2v/Yosys in the prebuilt toolchain so the ASIC flow's timing/power story works out-of-the-box from a fresh toolchain_install.sh --all.
  • Trap / exception foundation (preemption groundwork). Synchronous machine-mode trap path: ECALL / EBREAK / MRET with per-warp M-mode CSRs (mstatus, mtvec, mscratch, mepc, mcause, mtval). The branch unit is the trap producer (cause rides the existing branch-control interface); the scheduler snapshots mepc/mcause and redirects the warp PC to mtvec, with mret returning through mepc — mirrored one-for-one in the SimX scheduler (per-warp raise_trap / mret, replacing the old kill-all-warps trap behavior). Host-side ELF loading + HTIF tohost monitoring (sim/common/elf_loader.cpp, sim/common/host_monitor.cpp) are wired into both simulators, which is what lets the upstream riscv-tests suite run natively on Vortex. Asynchronous preemption itself (preempt-from-CP, interrupts, supervisor mode, register-file/IPDOM save-restore) builds on this and is future work. See docs/designs/trap_and_exception_foundation.md.
  • Toolchain refresh. LLVM 20, POCL 7.0, chipStar.
  • Kernel-entry calling convention (no callee-saved spills). vortex.kernel entries (__kernel macro, vx_spawn2.h) skip s0–s11/fs0–fs11 spills via an empty LLVM-20 callee-saved set (ra still saved); the KMU/__vx_cta_entry trampoline runs each once then vx_tmc zero. Why: generic-ABI per-CTA spills thrash the 16 KB L1 (sgemm2 +27% cycles).
  • Versioned toolchain pipeline for CI. VERSION is the single source of truth (VORTEX_VERSION, TOOLCHAIN_REV, GEM5_REV); CI cache keys + installer scripts both honour it; bumping a pin rolls the CI cache.
  • Perfetto trace integration. ci/perfetto.py renders RTL and SimX traces into Chrome Trace JSON (auto-detects flavour); see docs/perfetto_analysis.md.
  • AI-agent integration via AGENTS.md. Canonical entry point for AI agents and human contributors — foundation rules, documentation map, build/test/design invariants.
  • Test groups. New --vulkan, --gem5, --hip, --amo, --tensor_sp, --tensor_wg, --dtm, --mpi, --rvc, --vm, --graphics runners in ci/regression.sh; matrix expanded to rv32 + rv64.
  • Performance-regression gate (perf_gate). Catalog cases marked check: perf_gate run on rtlsim and gate against goldens in ci/baselines/perf/: retired instructions must match the baseline exactly (a mismatch means the workload changed, not performance), cycles must stay within ±2%, and a config hash invalidates a baseline whose run configuration drifted. The ratchet is two-sided — an improvement beyond 2% also fails until a human locks it in with pytest ci --update-baselines — so the goldens always state the real number. Coverage spans core, dxa, graphics, raytracing, and the tensor/tensor_sp/tensor_wg/tensor_mx families. Why: cycle regressions surface at the change that caused them instead of in a later benchmark sweep.
  • SimX↔RTL model-parity gate (model_parity). Catalog cases marked check: model_parity run the same app/args/configs on both drivers and assert (1) identical retired-instruction counts — a mismatch is functional divergence, not a timing gap — and (2) a cycle gap within the case's tolerance. SimX is the RTL's timing model; this check keeps the two in lockstep. Known modelling gaps carry an explicit known_issue: reason in the catalog rather than silently passing.
  • FPGA synthesis-regression gate (fpga_gate). ci/fpga_gate.py puts nine DUTs (cache, core, top, dxa, tcu, rtu, raster, om, tex) through Vivado synthesis + place-and-route at their target clocks and gates post-implementation Fmax and area (LUT/FF/LUTRAM/BRAM/URAM/DSP) against goldens in ci/baselines/synthesis/xilinx/ (default ±5%, per-build/per-metric overrides in ci/testcases/fpga_gate.yaml). Longest-first 2-up scheduling, per-build resume-from-checkpoint stamps, and an RTL-elaboration watch that fails a config typo in minutes instead of hours. The fpga tier is opt-in: it runs only on the self-hosted Vivado runner (.github/workflows/fpga_gate.yml) or by hand, never on hosted CI; goldens are re-recorded only by a human via --update-baseline. Why: the same golden/threshold/ratchet discipline as perf_gate, applied to what synthesis costs (timing closure, area) rather than what execution costs (cycles).
  • ASIC synthesis-regression gate (asic_gate). The open-source counterpart of the FPGA gate: ci/asic_gate.py puts the same hw/unittest DUT wrappers through sv2v → Yosys → ABC → OpenSTA on ASAP7 (RVT, TT) at their target clocks and gates post-synthesis Fmax and standard-cell area against goldens in ci/baselines/synthesis/yosys/, recording WNS/TNS, sequential and SRAM area, cell count and power alongside. The gating machinery is now shared: ci/synth_gate.py holds the catalog, config_hash, thresholds, known_issue, resumable sessions, scheduler and report, and fpga_gate.py/asic_gate.py are entry points that pin a tool adapter — fpga_gate behaviour is unchanged, config hashes included. DUTs are declared once per flow in hw/syn/{xilinx,yosys}/dut/catalog.mk, so make -C hw/syn/yosys/dut om and the gate run the same thing. The asic tier is opt-in and the gate runs from .github/workflows/asic_gate.yml — nightly, skipped when master has not moved, and fanned out to one hosted job per DUT because each takes 1-2 hours. Why: the flow needs no licence and no dedicated machine, so the discipline the fpga_gate applies on a self-hosted Vivado box can cover every contributor's changes. Two flows over the same modules also make a divergence between them a finding rather than noise. Three flow fixes fell out of building it. run_sta.tcl now reports report_worst_slack (signed) instead of deriving Fmax from report_wns, which clamps at zero and made every design that closed report exactly its target clock. configure now copies every *.mk fragment into a build tree, not just common.mk, so an included catalog.mk/build.mk resolves from build/ as well as from the source tree. And each synthesis flow's mandatory defines are now override CONFIGS +=: a CONFIGS= on a make command line silently defeats every CONFIGS += in a makefile, so the DUT dispatcher was dropping -DSYNTHESIS -DASIC -DYOSYS — the cache DUT synthesized to 1.08 M cells of flip-flop-inferred RAM with zero SRAM area. The Xilinx and Altera DUT flows carried the same trap (their gate passes CONFIGS through the environment, which appends rather than overrides, which is why it never fired there) and are hardened the same way; a command-line CONFIGS is now extra defines on top of the flow's, which is how it reads.
  • tests/opencl: full Rodinia suite + atomics benchmarks. The OpenCL test tree now carries the Rodinia benchmark suite wired into CI — backprop, bfs, b+tree, cfd, dwt2d, gaussian, heartwall, hotspot, hotspot3D, hybridsort, kmeans, lavaMD, lud, myocyte, nn (nearn), nw, pathfinder, srad, and streamcluster — plus atomics benchmarks (atomicreduce, histogram, in both tests/opencl/ and tests/hip/) that exercise the A-extension end-to-end across the L1/L2/L3 hierarchy, and the OpenCL image benchmark set (image_*; see the OpenCL image-support entry above). Why: moves OpenCL coverage from a handful of kernels to a standard, externally comparable workload suite that stresses the caches, atomics, and the image path together.
  • Hardware virtual memory (MMU v2). RTL translation hierarchy (hw/rtl/vm/): per-L1-cache socket TLB stage (non-blocking dTLB/iTLB + miss station), shared cluster L2 TLB (512e 4-way + megapage array), 2-walker PTW with walk cache; DCR-sourced satp, kill+report faults, page-granular allocation; vm CI on simx+rtlsim at XLEN 32/64. See docs/designs/virtual_memory_subsystem.md.
  • Multi-threaded SimX. Per-socket execution domains on lockstep worker threads (-DSIMX_MT=<T> via build CONFIGS), with cycles bit-identical across every thread count. New port-only communication API: SimEventLink control-plane links, KMU push dispatch matching the RTL launch bus, and framework-owned quiescence via SimPlatform::idle().
  • SimX functional emulation mode. -DSIMX_FUNCTIONAL (via build CONFIGS) selects a functional-only simulation kernel (unit latency, no backpressure, MT executor retained) behind the unchanged component API for full-speed architectural runs such as CTS conformance sweeps. Functional cycle counts are non-physical and excluded from perf_gate/model_parity.
  • CONTRIBUTING.md and this changelog at the repo root.

Changed

  • ALU, LSU and FPU scale with the issue width. VX_CFG_NUM_{ALU,LSU,FPU}_BLOCKS now default to VX_CFG_ISSUE_WIDTH (VX_config.toml); VX_CFG_NUM_SFU_BLOCKS stays at 1.
  • Shared memory scales with the lane count. VX_CFG_LMEM_LOG_SIZE is derived as VX_CFG_LMEM_NUM_BANKS x 2KB instead of a fixed 16KB, giving 8/16/32/64KB at NT=4/8/16/32 with bank depth constant.
  • DXA LMEM writes are counted and committed per banked row. The SimX drain counted one write per port request instead of one per row beat, and could leave a row half-written when the second port was full; lmem_writes now matches rtlsim exactly.
  • LMEM window decode no longer forms an end address. VX_lsu_slice.sv and VX_tcu_agu.sv test membership by comparing the address bits above the window, which stays correct when the window ends at the top of the address space.
  • TCU register-file bank-conflict-free mapping. TCU micro-op generation in TcuUopGen (and the SimX sim/simx/tcu/tcu_unit.cpp:1332 bank-conflict-free formulas) permutes A / B / C operand offsets so every uop's three RF reads land in different GPR banks; separate formula classes cover sparse, dense NT∈{4,16,64}, and dense NT∈{8,32}. Why: drops issue-stage stall cycles to zero on the TCU MMA loop — v2.x used a naive (step % sub_blocks) * block_size offset that incurred bank collisions on every other uop.
  • Profiling counters are host-driven via DCR (device-side dump removed). Kernel startup (vx_start.S) no longer dumps perf counters at exit; the host reads them on demand via vx_mpm_query → vx_dcr_read, gated by VORTEX_PROFILING. Why: drops perf-dump code from every kernel binary.
  • Source-tree consolidation. kernel/ + runtime/ → sw/{kernel,runtime,common}; new sw/common/ holds code shared by device and host (rvfloats, softfloat_ext, mem_alloc, plus the h/w-internal gfx_render.{h,cpp} host hardware model used only by simx). Why: removes include-path duplication and mirrors the sim/{simx,common,...} layout. Note: ABI types and downstream-visible config (tensor_cfg.h, the graphics on-wire types) live in sw/kernel/include/ so they ship through the install tree; sw/common/ is vortex-internal and never installed.
  • Downstream tools consume vortex via $VORTEX_PATH + pkg-config. mesa-vortex, pocl-vortex and chipStar no longer reach into the vortex source tree ($VORTEX_HOME) or build tree ($VORTEX_BUILD_DIR) for headers or libraries. They link against the install tree produced by make install through vortex-runtime.pc / vortex-kernel.pc. Mesa's vortex-runtime option is now vortex-path; pocl's -DVORTEX_PREFIX= / -DVORTEX_BUILD_DIR= are replaced by pkg_check_modules(VORTEX REQUIRED vortex-runtime) + -DVORTEX_PATH_{32,64} for per-XLEN device-side bitcode builds.
  • gfxutil.{h,cpp} removed. Binning() moved into graphics.h + graphics.cpp (now part of libvortex.so); the toVX{Format,Compare,StencilOp,BlendFunc} helpers (CGLTrace-specific test-input translators) inlined into tests/regression/gfx_draw3d/main.cpp, the only consumer; ResolveFilePath (test-asset filesystem resolver) inlined per-test, matching the resolve_path pattern raycast already used. Why: gfxutil.h was leaking cocogfx (CGLTrace, ePixelFormat) into what should have been vortex's public surface; eliminating it lets graphics.h be self-contained as the install tree requires.
  • riscv-tests .bin blobs removed from the source tree. Pre-built tests/riscv/isa/*.bin (c56562fe) and tests/riscv/benchmarks_{32,64}/*.bin (cd1656cf) collapsed into a single on-demand build under tests/riscv/common.mk that clones upstream riscv-tests at a pinned commit and builds per-XLEN behind a stamp file. Why: drops binary blobs from version control and pins behaviour to a single upstream commit instead of stale checked-in artifacts.
  • Configuration moved to TOML. VX_config.toml + VX_types.toml replace hw/rtl/VX_config.vh; ci/gen_config.py emits per-target headers (build/hw/VX_config.vh, build/sw/VX_config.h) and -D overrides from one source, with expr: / [[enum]] / [[builtin]] / [[param]] semantics. Why: gives the config typed scalars / cross-key expressions / typed enums the flat \defineblock could not express, and decouples SimX/runtime from the RTL source tree (no more-I$(ROOT_DIR)/hw`).
  • Configuration namespace. All HW config macros now carry the VX_CFG_ prefix; HW/SW layering split — VX_config.h is HW/sim-private. Why: resolves preprocessor collisions with LLVM/Clang and Verilator/SV \define`s, and stops HW config from leaking into kernel-side toolchain invocations.
  • No more global toolchain_env.sh. Each build/ carries its own resolved tool paths. Why: enables parallel multi-version Vortex worktrees on the same shell (the old source ci/toolchain_env.sh hijacked $PATH).
  • Build system. Tool-path env vars unified on the _PATH suffix; shared hw/syn/common.mk; standardized OPT_LEVEL across synthesis backends; LLVM default target fixed; dead hw/config call dropped. Why: normalizes the build/synthesis surface so cross-backend changes touch one place.
  • Verilator pinned to 5.028 (was tracking the latest 5.046 release in the prebuilt). Why: 5.046 removed the --xml-only flag that sim/opaesim/Makefile needs for --scope, tightened DEFOVERRIDE into a hard error breaking DEBUG=3 + -DGPR_RESET debug builds, and rejected cvfpu's C_PC / C_EXP_* macros under DPI_DISABLE + FPU_FPNEW (--config2). 5.028 keeps all three working (commit f00bb142).
  • SimObject channels — explicit fixed-size handshaking. Unbounded std::queue + blocking push/pop replaced by fixed-capacity channels + non-blocking try_send / try_pop (RTL ready/valid analog); [[nodiscard]] forces producers to handle backpressure at the call site. Why: makes SimX model true RTL backpressure (1:1 with ready/valid) so buffering bugs surface in C++ instead of only at RTL bring-up.
  • Scoreboard issue arbiter — GTO (Greedy-Then-Oldest). VX_scoreboard now drives a dedicated VX_gto_arbiter with a suppress mask for warps whose target functional unit is full (they keep aging but are skipped for selection until the FU drains). v2.x used VX_stream_arb (round-robin). Why: GTO matches mainstream GPU warp-issue policy — finish the currently-running warp before switching — yielding better ILP/cache locality than naive RR.
  • Scoreboard — XREGS dependency tracking + per-FU lock for multi-uop macros. Scoreboard gains a second plane (inuse_xregs) tracking FPU special registers (fflags, frm); every instruction declares rd_xregs / wr_xregs masks in decode (VX_decode.sv:377-380). New fu_lock / fu_unlock instruction flags let a uop sequencer hold a functional unit across a multi-uop macro — set by VX_tcu_uops.sv:422 for WGMMA (fu_lock=first_uop, fu_unlock=last_uop) — and the scoreboard refuses to issue any other warp's uop to the locked FU until it retires. Why: eliminates v2.x's coarse FPU-boundary stall for fflags/frm WAW/RAW hazards, and keeps multi-uop macros atomic on one FU instance — without fu_lock, interleaving warps would corrupt shared FU state (WGMMA's per-block A registers + accumulator).
  • Warp scheduler — per-warp ibuffer-capacity gate. VX_scheduler tracks per-warp ibuf_full and computes schedule_warps = ready_warps & ~ibuf_full, with an all_ibuf_full ? ready_warps : preferred_warps fallback to keep pipelines absorbing transient stalls. v2.x scheduled solely on active_warps & ~stalled_warps. Why: prevents the scheduler from issuing a warp whose decoded uops will only block on a full ibuffer downstream — wasting fetch/decode bandwidth and pushing back-pressure into the front end.
  • __syncthreads() (BAR opcode) now drains LSU before suspending warps so SMEM writes commit before any post-barrier reads. Gated on lsu_sched_drained in VX_wctl_unit.sv and lsu_drained() in sim/simx/sfu_unit.cpp.
  • BUG FIX: SimX cache-flush parity with RTL VX_dcr_flush — ProcessorImpl::flush_caches() (sim/simx/processor.cpp) now fans out to icache + dcache + {tcache, rcache, ocache} L1 surfaces in parallel (was dcache-only), matching the RTL shared-req/AND-of-done topology in VX_core.sv + VX_graphics.sv.
  • Sectored L2/L3 + word-size dcache banking. L2/L3 lines doubled to 2*MEM_BLOCK with a MEM_BLOCK sector (VX_CFG_{L2,L3}_SECTOR_SIZE) — one tag spans the line while fills/evictions transact at the bus granule, halving LLC tags. Dcache banks now derive from a reduced coalescer WORD_SIZE (NUM_BANKS = NUM_REQS = footprint/word), scaling with thread count. Why: cuts LLC tag area and lets dcache memory-level parallelism track NT without shrinking the line.

Fixed

  • SimX arbiter input grouping. TxArbiter/TxRxArbiter grouped inputs by log2ceil(inputs / outputs) where VX_stream_arb uses ceiling division, so a request count that is not a multiple of the output count left the trailing inputs unserved (3 cores over 2 shared icaches hangs at boot). Covered by the cache: demo-shared-icache-odd catalog case.