Conversation
… refactor Migrated from the gfx_v2 working tree (now retired). Adds the device-side software output-merger — depth/stencil test, blend, blend-modes, logic-op, write-mask — as the §6.5 always-correct fallback, plus the per-fragment ops as the single source of truth shared with the host FF model (§7): gfx_render.cpp's DepthTencil/Blender now forward to gfx_sw:: (no duplicated math). The ISA-agnostic per-fragment math validated bit-for-bit vs the v1 OM unit in the gfx_v2 tree (11 configs: depth read/write, blend over/add/min/max, alpha-saturate, write-masks). Compiles clean here. Runtime re-validation through the simx OM model is deferred until the in-flight OM v2 (vx_om4) migration settles — gfx_pipeline_om currently aborts independent of this change. A vx_om4-based device-SW-OM test is a follow-up once OM v2 lands. Note: large inlined SW merges need -mllvm -vortex-divergence-max-bbs raised (the divergence pass's default 100-BB guard); see gfx_sw::om_fragment. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…imx-validated) Fixes gfx rendering on simx, which was all-black at HEAD. Commit 4709fff moved the RASTER core + device FrontEndPool to the gfx_v2 §6.3 coarse-128px bin layout (rast_bin_header_t: dense headers + absolute-indexed sorted-pid array) but left host graphics::Binning() emitting the old rast_tile_header_t (8B, 32px tiles, interleaved pids). Every host-Binning consumer fed the new RASTER core mis-typed records -> zero covered quads -> all-black. Changes: - sw/runtime/graphics.cpp: Binning() now emits the coarse-bin layout the RASTER front end reads; gfx_draw3d bins at VX_CFG_RASTER_BIN_LOGSIZE. - tests/graphics/gfx_pipeline_{raster,tex,om}: oracle parse updated to the bin format (host Binning == device front end). - OM vx_om4 (sole windowed quad OM op): decode + SFU quad sequencer + OM unit (+ RTL), with the four OM-using gfx kernels converted from vx_om to vx_om4 staging the payload through the shared graphics window. - sw/runtime/graphics.h/.cpp: add FF register emitters program_raster/om/tex + draw() (genxml/si_emit-style state->DCR packing) for the driver/tests. - box_ref_128.png re-blessed for TEX v2 integer-mip/trilinear rounding (<=2 LSB; the reference predated TEX v2 by two weeks). Validated on simx: gfx_draw3d (triangle + box), gfx_raster, gfx_tex, gfx_om, gfx_pipeline_raster, gfx_pipeline_om all PASS (6/7). Pending (separate, pre-existing, not from this change): - gfx_pipeline_tex §8 one-batch-vs-two-batch multi-draw depth coherency (146 B edge-localized) in the full-residency batched-draw path. - RTL OM v2 (vx_om4) awaits rtlsim validation. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…p latency Two SimX RTU timing-model gaps vs the RTL, found by auditing which VX_CFG_RTU_* params SimX honors: 1. Context pool size. SimX sized its §8.6 async ray-context SlotPool by VX_CFG_RTU_CONTEXT_POOL (=32) — a param the RTL never references. The RTL sizes its in-flight context pool by VX_CFG_RTU_NUM_CTX (VX_rtu_core.sv, default = NUM_THREADS = 4). SimX was therefore modeling 8x the in-flight contexts → far less pool backpressure than the hardware. Now sized by NUM_CTX (with the same NUM_THREADS fallback the RTL uses). The allocate() path already stalls gracefully on a full pool, so the smaller pool only adds backpressure, never deadlocks. 2. Per-ray setup latency. The SimX cost model charged only BoxPe/TriPe cycles; the reciprocal-setup span the RTL scheduler waits before traversal (SETUP_LAT = RTU_FDIV_LAT = 17, the 1/dir pipeline depth) was free. Now charged once per ray (latched via Slot::setup_charged so callback-resumed segments don't re-charge it). Findings (no SimX change needed): - VX_CFG_RTU_RECIP_DSP_SEED is an area/resource knob, NOT timing: both recip backends honor the same SETUP_LAT span (VX_rtu_scheduler.sv:372). - VX_CFG_RTU_XFORM_LATENCY is unused by the RTL; the real per-instance transform latency is 4 * RTU_LATENCY_FMA = 36. Validated: all RTU tests pass on simx (incl. pool-stressing async_batch, reform_mw, recursive). Remaining for full <5% gap validation: per-instance transform latency (36, needs a SimX transform counter); STACK_DEPTH spill-stall modeling; and a RTU-traversal-dominated workload runnable on BOTH backends for cycle-gap measurement (rt_raycast dumps perf but is rtlsim-deferred by an RTL scoreboard deadlock under sustained multi-warp load). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The Vulkan draw3d test handled only FORMAT_A8R8G8B8 textures: it seeded
the staging buffer to 0xff and swizzled into it only for A8R8G8B8 with a
>=w*h*4 size. box.cgltrace stores its texture as R5G6B5 (16bpp), so the
size check failed, the swizzle was skipped, and a white block was
uploaded -- box rendered solid white through vortexpipe (the driver
faithfully uploads what it is given; the vx_tex4 path itself is correct).
Two fidelity fixes so the test matches the native draw3d reference the
PNGs were blessed against:
- Expand R5G6B5 -> R8G8B8A8 with the SAME bit-replication the Vortex
TEX unit applies (sw/common gfx_render Unpack8888), so the uploaded
texels equal the HW-expanded ones. A8R8G8B8 keeps its direct swizzle.
- Bind a bilinear sampler when the draw call's trace magfilter is not
NEAREST (the native path programs VX_DCR_TEX_FILTER the same way);
the test previously hardcoded VK_FILTER_NEAREST for every draw.
box now renders its stone texture correctly: pixel delta vs the native
reference drops from 3403 (all-white) to 17 (0.1%), the residual being
sub-pixel rasterization coverage / interpolation precision between the
vortexpipe and native pipelines at texture-detail boundaries.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Two v2.1 P0 items from the RTU reviews (rtu_runtime P0-2, rtu_simx P0-3):
- Scene header word2 (raytrace.h build_bvh_scene) wrote 0 where the
VxBvhSceneHeader contract puts `scene_bytes` (rtu_memory.cpp:127-130
reads hdr+8 to size the RtuCore pre-fetch). With 0 the engine fell back
to pre-fetching the whole per-lane budget (kRtuMaxBvhSceneBytes) every
ray. Write the real serialized size so the pre-fetch is exact; word3
stays leaf_count (the comment wrongly called word2 node_count).
- The SimX cost model never charged the TLAS object-space transform.
Capture the bvh_instance_descents delta exactly like the box/tri PE
deltas in compute_intersections and add kRtuXformLatency (= 36 =
4*RTU_LATENCY_FMA, VX_rtu_xform.sv:62) per descent. Promotes the
"(Pending)" note in rtu_types.h to a real constant.
Validated: rtu_smoke_bvh_instanced (2 instances of 1 BLAS) PASSED on simx.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Three v2.1 P0 contract/ABI items:
- pid width (gfx_runtime P0-1): the bin-sort key packs the pid into its low
PIPE_PRIM_BITS (=20) but the RASTER unit addresses it with a
VX_RASTER_PID_BITS (=16) field, so a scene with >65536 visible primitives
silently aliased on the device. Add a static_assert(PIPE_PRIM_BITS >=
VX_RASTER_PID_BITS) plus a hard runtime reject + diagnostic in Binning()
and FrontEndPool::append (the 16-bit consumer width is the real cap).
- f0-f7 ray window (rtu_kernel P0-2): the trace2 decoder reads the ray from
f0..f7 by HW convention with no runtime check — the largest silent-failure
surface. Add _Static_asserts that vx_ray_t is exactly the 8 contiguous
floats {origin,dir,tmin,tmax} those register-pinned asm operands map onto,
so a field reorder/padding change fails to compile.
- config-lane table (rtu_kernel P0-1): update proposal §5.1 + the §2 taxonomy
to the SHIPPED 3-lane wgather layout (lane0 = suppressed self slot, lane1 =
scene, lane2 = payload, lane3 = flags[15:0]|cull[31:16]); the table is
normative for the RTL and Mesa lowering.
Validated: rtu_smoke PASSED (f0-f7 asserts compile); gfx_pipeline_tex
device==Binning + host-Binning image checks PASS (pid guard compiles, no
binning regression).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The BVH walker capped its stack at a fixed 16 and silently dropped pushes
on overflow (rtu_walker.cpp), so a deep sub-tree could be skipped entirely
— a missed-closest-hit correctness divergence from the HW, which trail-
restarts to re-find evicted subtrees and so still visits every leaf
(v2.1 RTU-simx P0/consol P0-6).
SimX is the correctness oracle, so make the functional stack unbounded —
it never misses a hit. The HW short-stack depth (VX_CFG_RTU_STACK_DEPTH)
overflow is instead counted as a trail RESTART (new bvh_stack_restarts
perf counter) and charged in the cost model as ~VX_CFG_RTU_STACK_DEPTH
re-descended box-PE issues per restart (approximate, pending the full
trail model). Shallow scenes are unchanged (no overflow → no restart).
Validated: rtu_smoke_bvh_{basic,multilevel,instanced} + cull_back PASSED.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The SimX BoxPe/TriPe cost model assumed a W-wide parallel PE array
(÷VX_CFG_RTU_BOX_PE=4) at latency 4/6 — but the RTL has ONE box PE and ONE
tri PE per RtuCore, each streaming one primitive per cycle, with much deeper
pipelines. The result was a ~10x throughput overstatement (the headline
SimX<->RTL timing-parity defect, v2.1 RTU-simx P0-1/2 / consol P1-1) and
BOX_PE/TRI_PE/NODE_LATENCY/TRI_LATENCY are dead config in the RTL.
Re-base cycles_for on the actual RTL depths, expressed symbolically so they
track the FP latencies:
- box = 1 issue/box + (3*FMA + 4) = 31-cycle drain (VX_rtu_box_pe.sv:76)
- tri = 1 issue/tri + (8*FMA + FDIV + 2) = 91-cycle (VX_rtu_tri_pe.sv:68)
with kRtuLatencyFma=9, kRtuFdivLat=17. Drops the ÷4 the HW doesn't have.
Validated: rtu_smoke{,_bvh_multilevel,_ahs,_chs} PASSED (timing-only change,
hits unchanged).
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
OmCore runs up to kInflight slots concurrently and drain_req_in admitted any request into a free slot with no check that an in-flight slot targets the same pixel. Two same-(x,y) fragments could both READ stale depth/colour before either WROTE, so the last write won by slot scheduling rather than submit order — a latent overdraw-ordering inaccuracy vs a real ROP, which serialises same-pixel fragments (v2.1 gfx-simx P0/P1, separate from the §8 multi-draw issue). Hold a request whose covered pixels collide with any in-flight slot until that slot retires; other pixels on other channels still progress. Mirrors the RTL ROP per-pixel ordering. Validated: gfx_pipeline_tex device==Binning + host-Binning image checks PASS (no regression). §8 multi-draw diff unchanged at 146 bytes, confirming (as the review found) this is independent of §8. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
… oracle The host/device cull-clip parity hazard (v2.1 gfx-runtime P0-2): the device front end applies SETUP_CULL_* back-face culling and near-plane clipping (pipe_frontend.h pipe_clip_and_setup), but host Binning() back-face *flips* winding and skips that clip — so the "device == Binning()" cross-checks were only incidentally bit-exact (the test scenes are front-facing and within the near plane). Per the review's sanctioned option, demote Binning() from "bit-exact oracle" to "coverage reference": document in graphics.h that it does not replicate the device cull/clip and that the device is authoritative for culled/clipped scenes, and annotate the gfx_pipeline_tex cross-check so its bit-exactness is understood as scene-conditional, not a general guarantee. Also fixes the stale rast_tile_header_t doc (now the §6.3 coarse-bin layout). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
build_bvh_scene<W> emitted a single leaf (root IS the leaf holding every triangle) — a valid but unpartitioned BVH that never exercised the walker's internal-node datapath (box tests, ordered descent, sub-tree pruning), the whole reason the RTU exists. Replace it with a real top-down spatial builder (detail::BvhBuilder): - Binned-SAH binary partition (12 bins, surface-area cost), applied repeatedly to fan out up to W children per internal node; index-median fallback for coincident centroids guarantees progress. - Post-order serialization so children precede parents and every child_offset is absolute from the scene base; root is the last node. - Conservative AABB quantization about a per-node origin + per-axis int8 exponent (min rounds down, max rounds up) so the walker's reconstructed child box always contains the true box — never culls a real hit. Leaves hold ONE triangle. The leaf format reports gl_PrimitiveID as prim_base + within_leaf_index, so a multi-tri leaf can only preserve Vulkan primitive IDs for a contiguous ascending source-index run, which spatial SAH does not produce. One tri per leaf sets prim_base to the source index, so the reported primitive ID is always exact. Multi-tri leaf clustering needs a per-leaf prim-ID remap field (a format extension) and is the documented follow-up. Internal-node byte-layout constants move to the shared rtu_cfg.h (the host cannot include the SimX-only rtu_bvh.h); the host refuses to emit a scene larger than the walker's pre-fetch budget (RTU_BVH_MAX_SCENE_BYTES). New test rtu_bvh_multinode: 36 spatially-distributed triangles -> asserts the emitted tree has internal nodes (21) with all 36 prim IDs preserved once, then fires one ray per triangle and matches a host brute-force Moller-Trumbore oracle (status/prim/t/u/v) through real multi-level traversal. simx-only on rtlsim (sustained-load profile, same as rt_raycast). Validated on simx: rtu_bvh_multinode PASS, rtu_smoke_host_cfg PASS (single-leaf root path preserved), rtu_smoke_bvh_basic / rtu_smoke_bvh_instanced PASS. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The runtime (sw/runtime/graphics.cpp) and the SimX raster oracle (sim/simx/raster/raster_core.cpp) emit/consume the gfx_v2 §6.3 coarse-bin layout: a dense block of 12-byte rast_bin_header_t records (bin_x/bin_y + 32-bit absolute pids_offset/pids_count) followed by one shared sorted-pid array, with bins at 128 px (BIN_LOGSIZE). The RTL front end still read the old 8-byte rast_tile_header_t (16-bit fields, header-relative pids_offset, 32 px tiles) — a format mismatch that left the RTL raster path unrenderable on the new buffers. VX_raster_mem.sv: - TILE_HEADER_SIZEW 2 -> 3 words; fetch 3 header words (mask 9'b111) at both header-fetch sites (initial + next-tile refetch). - Decode pids_offset/count as full 32-bit words (word1/word2) instead of the packed 16-bit halves of word1. - pids_addr is now absolute: pids_array_base = tbuf + tile_count*3 words (the dense header block shared across the N striped instances) + pids_offset, replacing the old header-relative "mem_req_addr[1] + offset + 1". VX_graphics.sv: - Instantiate VX_raster_core with TILE_LOGSIZE = BIN_LOGSIZE so the front end's top-level "tile" is the 128 px bin and VX_raster_te refines bin -> block -> quad, matching the runtime bin granularity and the SimX oracle's bin_x*bin_size origin. (te TILE_FIFO_DEPTH grows 4^(BIN-BLOCK); its BRAM / 300 MHz cost is a U55C synthesis-pass item.) Validated on rtlsim: gfx_raster output matches the SimX oracle reference bit-for-bit (CompareImages 0 errors). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The per-(warp,lane,slot) graphics register window — the slot file that SETW writes and vx_tex4 / vx_om4 / the RTU stream their operands through — was owned by RtuUnit and gated on VX_CFG_EXT_RTU_ENABLE. So a build with OM and/or TEX but no RTU could not source the window at all: SETW failed to decode (the gfx pipeline tests had to co-enable RTU purely to get the window), and the SimX vx_om4 path had a degraded no-RTU fallback. This is the v2.1 P0-3 cross-component coupling. The RTL already decouples it (VX_gfx_window_pkg.sv lives at the SFU level, not in the RTU core); SimX now matches. Extract a standalone GfxWindow (sim/simx/gfx_window.h): the [warp][lane][slot] storage plus get/set and the pure window ops process_set (SETW) and process_getw_uop (GETW/GETWF). SfuUnit owns it under a new derived gate VX_GFX_WINDOW_ENABLE (= OM || TEX || RTU); RtuUnit borrows it by reference for its trace/wait/terminal streaming. - types.h: RtuType / IntrRtuArgs and their OpType / IntrArgs variant membership move from VX_CFG_EXT_RTU_ENABLE to VX_GFX_WINDOW_ENABLE. SETW is a shared window op; CB_RET / TRACE2 / WAIT2 stay RTU-only ops (still only decoded with the RTU built). GETWF / GETW remain RTU-gated in decode (macro-ops expanded by the RTU sequencer, only emitted by RTU dispatchers / vx_rt_wait2). - decode.cpp: split funct3=6 — SETW under VX_GFX_WINDOW_ENABLE, CB_RET/GETWF/GETW under VX_CFG_EXT_RTU_ENABLE; widen the RtuType disassembler visitor. - sfu_unit: own GfxWindow; route vx_tex4 / vx_om4 window reads + writes and the SETW/GETW dispatch to it; the RTU-only CB_RET/TRACE2/WAIT2 handlers stay RTU-gated. Drop the now-dead vx_om4 no-RTU fallback. - rtu_unit: borrow GfxWindow&; drop the owned regfile_ + process_set / process_getw_uop (moved to GfxWindow). The three gfx pipeline tests (gfx_pipeline_om / gfx_pipeline_tex / gfx_draw3d) no longer co-enable RTU; they now exercise the OM+TEX-without-RTU path directly. Validated on simx: - RTU on: rtu_smoke + rtu_bvh_multinode PASS (no regression). - RTU OFF (RASTER+OM+TEX only): gfx_pipeline_om RESULT PASS (image bit-exact), gfx_draw3d PASSED (full FF pipeline), gfx_pipeline_tex FF checks PASS (device==Binning + image vs host-Binning); its only failure is the separate, pre-existing A1 §8 multi-draw diff, unrelated to the window decouple. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ite commit) A depth-tested multi-draw frame (near quad then far full-screen quad into a shared depth attachment, DEPTH_FUNC_LESS) rendered DIFFERENTLY as one CP batch vs two host batches — a 146-byte (~36-pixel) divergence at the near/far overlap. Root cause: the OM same-pixel R-M-W interlock (collides_with_inflight) holds a candidate fragment only while a same-pixel slot is in_use, but a slot reached DONE and freed at depth-write ISSUE — before the write committed. The ocache is write-through with NO write response, so writes were fire-and-forget. When the two draws overlapped in the OM pipeline, the far draw's depth READ could issue right after the near draw's slot freed but before its depth WRITE committed, reading stale (cleared) depth → the far quad wrongly passed DEPTH_FUNC_LESS and overwrote the near quad. Two separate batches drained fully between draws, so they were coherent — hence the one-batch vs two-batch mismatch. A real ROP orders same-pixel fragments through write COMMIT, not write issue. Match that: - cluster.cpp: enable ocache write responses (the OM is its only, auxiliary consumer). - om_core.cpp: register each write in pending_mem_ (is_write) so its commit ack is tracked; drain_mem_rsp marks z_write_acked / c_write_acked; advance_write_ issue holds the slot in WRITE_ISSUE until every write is BOTH issued AND acked before DONE. The slot (and its same-pixel interlock) now spans the full read-modify-write-COMMIT window, so a later same-pixel read can't bypass an in-flight write. Validated on simx (RTU off, RASTER+OM+TEX): - gfx_pipeline_tex: depth multi-draw one-batch == two-batch — RESULT PASS (was FAIL, 146-byte diff). - gfx_pipeline_om / gfx_draw3d / gfx_om / gfx_raster / gfx_tex: all PASS (no regression). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ping The RTU's in-flight ray-context pool is decoupled from SIMD width by VX_CFG_RTU_NUM_CTX (default = NUM_THREADS). The rtu_rtl_efficiency proposal flagged the NUM_CTX != NUM_LANES ray->context mapping as an unresolved / deferred design point, leaving open whether NUM_CTX > NUM_LANES was even functionally correct. Verified it IS correct: with NUM_CTX=8 > NUM_LANES=4 a warp's NUM_LANES rays occupy the low contexts and the extra contexts idle (VX_rtu_core.sv lines 73-75; SimX SlotPool sized by VX_CFG_RTU_NUM_CTX). rtu_smoke at NUM_CTX=8 produces the correct hit on BOTH simx and rtlsim. There is no correctness bug; the spare contexts are simply unused (the core accepts one warp end-to-end today). Exploiting them for concurrent multi-warp throughput is a separate optimization whose payoff is FPGA-measured and so lands with the synthesis pass. Lock the correctness in permanently: - New test tests/raytracing/rtu_smoke_numctx (reuses the rtu_smoke sources with -DVX_CFG_RTU_NUM_CTX=8) — registered in the suite, runs on simx AND rtlsim (not rtlsim-deferred: verified PASS on both). - Update rtu_rtl_efficiency.md §3 / open-items: the mapping is implemented + verified; concurrent multi-warp use is the remaining synthesis-pass throughput optimization. Validated: rtu_smoke_numctx PASS on simx and rtlsim. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The on-device setup+binning front-end ABI (expand_arg_t / setup_vertex_t / pipe_arg_t / SETUP_*) is the shared host/device contract a downstream driver launches expand_k/setup_k/binning_k against — its own charter names vortexpipe as a consumer — but it was never installed, so consumers' -I$VORTEX_PATH/ kernel/include couldn't resolve it. Install it alongside vx_gfx_abi.h, the same sw/common shared-ABI category. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The single-draw `depth` test renders both triangles in one vkCmdDraw, so it cannot tell whether the depth attachment survives BETWEEN draws. multidraw issues the near (blue, z=0.25) and far (red, z=0.75) triangles as two separate draws (two pipelines, not gl_VertexIndex/firstVertex — vortexpipe's self-contained VS ignores the vertex-index base) sharing one renderpass depth clear. With a persisting depth buffer the LESS test rejects the later far triangle at the centre, which stays blue; a driver that re-clears or loses depth between draws lets far win and the centre comes out red. This is the cross-draw-residency check the vortexpipe P1-6 work needs to validate against — it failed (centre red) before the depth-residency fix. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Add a NO_OM knob to the vulkan test common.mk (builds the device WITHOUT the OM
unit) and two tests that exercise vortexpipe's P1-7 SW output-merger fork — the
output-merger running in SIMT software when the OM unit is absent, never falling
back to the host:
- swom — the multidraw depth-persistence scene with no OM unit: proves
the SW OM's depth test + framebuffer write + cross-draw depth
persistence are correct.
- swom_draw3d — the draw3d scene with no OM unit: the rendered PNG must be
BINARY-EQUAL to the same golden reference the HW-OM path
matches, proving the SW OM is colour-correct (bit-exact), not
just depth-ordering correct.
Both reuse existing scenes/assets; only the device config (no VX_CFG_EXT_OM_
ENABLE) differs.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…+ dxa/opencl)
Brings in the 30 commits from the vortex_ci tree: CI 2.0 (catalog-driven,
driver-sliceable pytest architecture), scoreboard issue-stage timing close,
elastic cache bank pipeline + AMO timing fix, dispatch carry-save adder,
CTA-dispatch pipelining, vx_max_occupancy_grid 1D fix, K-major DXA DMA model,
OpenCL ICD-only / system-wide registration, RHEL-family configure detection,
and assorted cleanups — onto prism's graphics + RTU work.
Conflicts resolved:
- .github/workflows/ci.yml -> theirs: adopt the CI 2.0 dynamic test matrix
(include: fromJSON(plan.outputs.cells)). The catalog (ci/testcases/*.yaml,
incl. vulkan.yaml + graphics.yaml) drives coverage, so the prism vulkan
suite (multidraw/swom/swom_draw3d) is still exercised via run-{driver}.
- sw/runtime/common/vortex2_internal.h -> ours: keep prism's
cp_submit_launch_qmd (the CP QMD launch) + the fuller cache-flush comment.
Validated post-merge: regenerated the config headers from the merged
VX_config.toml (it added VX_CFG_L3_LATENCY for the elastic cache pipeline) and
the full vulkan graphics suite passes on the rebuilt SimX —
draw3d (bit-exact), depth, multidraw, swom (no-OM SW output-merger), triangle,
compute.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The RASTER producer kept a shadow `pending_count_` counter alongside the `pending_reads_` map and advanced its load FSM on `pending_count_ == 0`. The two could desync if a memory response tag were ever dropped (the counter would never reach zero), wedging the producer. Gate the LOAD_TILES/PIDS/PRIMS -> next transitions on `pending_reads_.empty()` directly and drop the redundant counter — one source of truth for outstanding reads. Validated: draw3d (bit-exact), triangle, depth, multidraw all pass on the rebuilt SimX. (gfx_simx P1.) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The bin-base stage ran the exclusive prefix-sum of bincount[] (and the dense rast_bin_header_t writes) on a single lane (tid==0) in an O(B) loop, serialising the whole CTA for large bin grids. Replace it with the same cooperative chunked-scan the BSCAN stage already uses: each thread sums its bin chunk into btsum[], lane 0 scans the per-thread partials, then every thread writes its chunk's binbase + headers from its scanned base offset; the non-empty-bin count (meta[2]) is a final btsum[] reduction. BBASE is single-CTA, so __syncthreads across the T lanes is valid, and btsum[] is free at this stage (BSCAN already consumed it). Bit-exact: the binning output is unchanged — gfx_pipeline_raster's device== Binning() cross-check PASSes at 128/256/512 (B = 1/4/16 bins) and draw3d stays pixel-exact vs the golden reference. (gfx_kernel P1-3, BBASE half.) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
om_fragment (the SIMT software output-merger) only compiles correctly when the fragment kernel is built with -mllvm -vortex-divergence-max-bbs: its full depth+blend+ROP merge overruns the Vortex divergence pass's default 100-BB guard, and on a trip the pass silently skips StructurizeCFG and miscompiles the kernel. That is a build flag, not a source change, so it can't be enforced in the header by itself. Encapsulate it: add sw/gfx/libgfx_sw.mk exporting LIBGFX_SW_VX_CFLAGS (the flag + -DGFX_SW_DIVERGENCE_OK), and guard om_fragment with `#error` when a device build pulls it in without GFX_SW_DIVERGENCE_OK. A miscompile becomes a loud build failure instead of silent breakage. No current device kernel includes gfx_sw.h, so nothing trips today; the host include path (gfx_render.cpp, __VORTEX__ undefined) is unaffected — draw3d rebuilds + passes. (gfx_kernel P1-2.) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
A quad's 4 corners (gl_VertexIndex-driven) rendered via vkCmdDrawIndexed with a UINT16 index buffer [0,1,2, 0,2,3] (two triangles sharing the 0-2 diagonal, vertices 0 and 2 reused). Validates the vortexpipe on-device indexed path: gl_VertexIndex resolved through the index buffer (a driver using the raw thread id would read past the 4 corners), the VS output written per index position, and the front end assembling the two triangles. Under MESA_VORTEX_STRICT=1 a PASS means the indexed draw ran on the Vortex device, not the llvmpipe fallback. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…item 6) The RASTER producer ran the entire tile-traversal walk to completion in a single tick and transitioned straight to READY, so the carefully order-accurate TE/BE walker contributed ZERO cycles — the RASTERIZE stage had no quad-emission rate and only the consumer side (one quad/cycle) modeled any throughput. te_walk_tile now returns its pipeline-cycle count (one per loop iteration: one block emit or one subdivision push — the walker is a 2-stage FIFO pipeline that advances once per cycle). run_rasterizer sums it into walk_cycles_, and the FSM charges that as a RASTER_DRAIN countdown before reaching READY. Because serve_consumers() runs only in READY, this is observationally identical to stepping the walker one stage per tick (warps stall until the producer has paid its emission latency) without restructuring the walk — the quad order/content is unchanged. Adds a raster_cycles perf counter for the walker rate. Validated on simx: gfx_pipeline_raster / gfx_pipeline_om / gfx_draw3d all PASS (functional output bit-identical; only the modeled cycle count changes).
…nt B item 7b) The setup front end staged every clipped/setup triangle into a per-triangle scratch slot (slot_prim, 120-byte rast_prim_t + slot_bbox) in SETUP, then EMIT read each survivor back and wrote it to the compacted dense primbuf — a full round-trip of the largest front-end record purely to remove the gaps left by culled sub-triangles. Fuse it: SETUP now only counts survivors (keep[t]); EMIT recomputes clip+setup and scatters each rast_prim_t straight into prim[offset[t]+s]. clip+setup is a pure function of the unchanged input vertices, so the SETUP count and the EMIT per-triangle output match exactly. This halves primbuf write traffic (one dense write instead of slot-write + dense-write) and drops the slot_prim/slot_bbox scratch buffers entirely (NT*MAX_SUB*120B + NT*MAX_SUB*16B of device memory). The only added cost is re-running the cheap front-end clip+setup math, which is negligible against the per-pixel fragment/OM stages. Consumer contract (dense prim[0..P] / bbox[0..P]) is unchanged, so RASTER/SimX/RTL are untouched. Validated on simx: gfx_pipeline_raster / gfx_pipeline_om buffer cross-check "device == Binning" PASS (bit-identical dense primbuf), gfx_draw3d golden PASS.
SETW / GETWF / GETW decode the window slot from funct7[6:2] (5 bits) but masked with &0x3F (6 bits). funct7>>2 is already ≤0x1F so this was behaviorally inert, but &0x1F now matches the RTL field width exactly and documents the 5-bit slot, removing a latent mismatch if the encoding ever widens.
…y matrix)
Phase 0 of the gfx_v2 master plan (docs/proposals/gfx_v2_true_gpu.md §1.3/§6):
mechanize the C1-C5 FF<->SIMT binding law so the recurring un-ordered-handoff
bug class fails the build or the matrix instead of slipping through review.
- sim/simx/gfx_doctrine.h: classify every FF op (CUSTOM1 -> SFU) at the decode
chokepoint into Scoreboarded / SideEffectFree / KnownViolation, checked against
its decoded destination. An unclassified FF op aborts unconditionally (a new op
cannot ship without declaring its handoff). vx_om4 and SETW are the enumerated
§3 KnownViolations: warn-once by default, fatal under VX_GFX_STRICT_DOCTRINE=1
(the switch P1/P2 flip on as each defect is retired).
- sim/simx/decode.cpp: invoke the assertion at the end of the EXT2 case.
- ci/testcases/graphics_parity.yaml: new graphics_parity category sweeping
multi-core x multi-cluster x multi-drawcall on both simx and rtlsim vs a shared
golden (agreement with one reference == SimX<->RTL parity). This is the missing
dimension that hid the multi-core handoff bugs.
- docs/proposals/gfx_v2_interface_law_checklist.md: the C1-C5 PR checklist.
Validated: assertion compiles in all configs, warns-once + test passes by
default, aborts under strict. Matrix runs (P0 gate, red allowed): box-{1,2,4}core,
box-4core-2raster, box-2clu-2core, box-1core(rtlsim) PASS; evilskull-2core fails
(449 px, OM blend-ordering) — the §3.1 OM-determinism defect P1 turns green.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…hread
P1 / gfx_v2 §3.1. The vx_om4 SFU dispatch caches the per-lane descriptor and
window base at sub-pixel 0 (om_desc_/om_base_), then streams the quad's four
sub-pixels across ticks — reusing trace->src_data in place to carry each
sub-pixel's pos/colour/depth to OmUnit. That overwrite destroys the desc/base
operands after sub-pixel 0. When the op is re-presented to the same SFU block
before it retires (observed under multi-warp × multi-thread SFU contention), the
F==0 capture runs a second time and re-reads the now-clobbered operands —
om_base_ becomes (colour & 0x1f), the window read lands on an empty slot, and the
fragment is written black. The result was a deterministic drop of one warp's
fragments at a draw edge (the gfx_om bottom-row failure; 4–12 px), invisible to
single-warp or single-thread runs.
Fix: capture desc/base exactly once per op, gated on a new per-block om_captured_
latch (set at first entry, cleared at retirement) instead of on F==0, so a
re-entry cannot re-read the overwritten operands. Minimal and self-contained;
vx_om4 stays fire-and-forget (the architectural scoreboard handle for the
FF↔SIMT window handoff remains the P3/RTL item — om4 is still a §3 KnownViolation
in gfx_doctrine.h).
Validated on SimX: gfx_om {plain,depth,blend,depth+blend} 4/4 PASS at the
failing 4-warp×4-thread config (were 4–12 px each); gfx_draw3d box (1/4 core) and
triangle, no regression. Localized via w1t1/w4t1/w1t4 PASS, only w4t4 fails.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…scenes
The parity matrix's evilskull cell compared against a committed golden that
neither SimX nor rtlsim matches byte-exactly (they agree with each other but are
±1 LSB off the reference — a texture/interpolation precision delta, not an
FF↔SIMT handoff bug). A parity-matrix cell must use a byte-exact golden so a red
cell signals a real cross-config divergence, not stale reference data. Replace
the evilskull cell with triangle (a distinct setup/coverage path) at 2 cores.
With the gfx_om OM window-cache fix (prev commit), the SimX half of the matrix is
now fully green: box {1,2,4}-core, box 4-core/2-raster, box 2-cluster/2-core
(multi-drawcall), and triangle 2-core all byte-exact.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Give each perf_gate benchmark a same-app/config model_parity twin so a perf-regression point is always backed by a simx<->rtl parity check: sgemmx/stencil3d/raycast (model_parity.yaml), sgemm2-fp16 (tensor.yaml), and the six DXA cases (copy-mcast, sgemm2-dxa[-mcast], wgmma-dxa, wgmma-sp-dxa, wgmma-dxa-mcast). fp16 TCU twins stay at the existing K=64 datapath-valid size. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
frag_param_r is XLEN wide but its only always-on consumer truncates it to VX_CFG_MEM_ADDR_WIDTH (48 on 64-bit builds), and the full-width read lives in a trace block that non-debug builds compile out. Verilator -Wall then fails the xrtsim driver build with UNUSEDSIGNAL on bits [63:48]. Tag the register with UNUSED_VAR, matching VX_kmu's dcr_param idiom. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Retarget perf_gate-wgmma-int8 to the fp16 datapath (perf_gate-wgmma-fp16) at the same K=128 depth, and regenerate the rtlsim baseline (758714 cycles / 525536 instrs). fp16 needs the sgemm_tcu_wg ULP gate raised from 10 to 16. The TCU chains K/tcK dot-product ops, each rounding its fp32 accumulator once, while matmul_cpu's reference accumulates in double with no intermediate rounding; the gap grows with K and measures 16 ULP at K=128. The RTL datapath is verified bit-exactly against its windowed reference by hw/unittest/tcu_fedp (run-tcu, fmt=2, 100k vectors), so this bounds a reference idealization rather than relaxing a hardware tolerance. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
The plru/fifo store read address only advanced on fill cycles, so the victim way for a fill was computed from the line captured at the previous fill — stale for any non-back-to-back fill. Use the tag-store idiom (read enable = ~stall, read-first) so the look-ahead always tracks the request entering S0. rtlsim perf_gate: tcu-fp16 -9.9%, sparse -10.4%, draw3d -12.6%, wgmma-dxa -9.8%, fp16-mc -7.8%; sgemv +14.8% and wgmma-int8 +11.5% regress from losing the accidental scan-resistance of stale victims (policy follow-up tracked separately); vecadd/sgemm-mc unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The command-list submission paths emitted CMD_LAUNCH_QMD unconditionally, but only the emulation CP decodes it — the RTL CP silently skipped the launch and, worse, re-parsed the descriptor payload bytes at the default command stride, executing them as phantom commands (garbage DCR writes and memory transfers). Every vulkan draw on the opae/xrt drivers lost its launches this way. Mirror the SUPPORTS_DRAW pattern: a new SUPPORTS_QMD DEV_CAPS bit (26) advertises the QMD decoder; the emulation CP sets it, the RTL CP does not. Without it the runtime replays the staged descriptor as plain CMD_DCR_WRITEs followed by CMD_LAUNCH — functionally identical, more ring commands — and skips the device-side QMD staging. OP_DRAW additionally requires the QMD bit since draw bundles embed LAUNCH_QMD steps. Harden the RTL CP unpacker against the same failure class: an unknown opcode now ends the cache line (fail-stop) instead of walking its payload as commands. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The opaesim/xrtsim Makefiles lacked the RTU package/include block the
rtlsim Makefile has, so any RTU-enabled config failed to elaborate on the
AFU drivers. Adding it exposed ~60 Verilator -Wall findings that only
rtlsim tolerated (it passes -Wno-fatal; the AFU sims correctly do not).
Resolve them at the source instead of spreading -Wno-fatal:
- tie the fmt pin ('0 = FP32) on the VX_fncp_unit comparator instances
- explicit width casts on intended truncations/expansions (stack index,
mem tags, prim index, ray-flag mask constants)
- hoist the ordered-insertion index (ins_pos) out of the sequential block
into combinational logic, removing the blocking-in-ff idiom
- size the round-robin candidate index instead of a 32-bit integer
- narrow u8_to_f32/pow2_f32 function temporaries to their consumed width
(bit-identical results; the fp32 leading one is implied)
- UNUSED_VAR tags for partially-consumed buses and the standard
IGNORE_UNUSED wrapper on VX_rtu_pkg
- static-assert the scheduler tag fits the rtcache tag field, which the
new cast would otherwise silently truncate on large-context configs
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…er engines The raster frame kick was inferred from a DCR decode: the engine self-started on a write to RASTER_FRAG_PARAM_HI, relying on a software write-order convention. The kick could fire before the launch's KMU startup DCRs were even submitted, so a fragment warp minted early fetched PC=0 (won by rtlsim's fast command processor, lost at the opae ring cadence), and nothing tied the frame to the launch fence that was supposed to cover it. The KMU is now the launch delegate: a kernel launch with an empty grid is a delegated draw launch by ABI definition. On such a start the KMU issues a payload-less kick token over a dedicated control interface (VX_raster_launch_if), distributed by VX_raster_launch_fork — a thin typed wrapper over VX_stream_fork in EAGER mode, whose delivered_r keeps the kick presented to each engine until that engine accepts and completes the master handshake only when all have. The KMU holds the device busy (raster_start_r) until the last engine fires and each engine holds it from its own fire, so the launch fence provably observes the frame — including empty frames that drain immediately. The engine arms on valid/ready fire (ready = ~armed_r); its drain logic is unchanged. FRAG_PARAM_HI reverts to plain configuration state. SimX mirrors the delegation: Kmu::launch_delegated() forwards the kick per cluster after the per-launch reset, the FRAG_PARAM_HI handler only latches the descriptor, and the arm no longer needs to survive reset. Runtime, mesa, and both command processors are unchanged — the grid-less FS launch and QMD staging already emit exactly this command stream. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Add the AGENTS.md testing rule that RTL cycle-moving changes must land together with the matching SimX timing-model update, enforced by the model_parity CI gate (exact retired-instruction match, cycle agreement within case tolerance; never widen a tolerance to absorb a divergence). Also tidy declaration alignment in VX_raster_dcr/VX_raster_dispatch. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Complete the read-prefix of a fill's pending chain directly from the
staged fill sector, bypassing the bank pipeline and the input arbiter;
the first write/AMO entry falls back to the replay path so coalesced
same-line accesses keep program order. The freed arbiter slots serve
demand requests while a chain drains. The SimX model mirrors both
effects plus the shared one-response-per-cycle port (pipeline
priority).
Forwarding removes the accidental serialization that kept core
requests out of the bank right after a fill, exposing two latent
races fixed here:
- tag store: the read-first bypass was keyed on the request's tag, so
a request to the line evicted by the previous cycle's fill read
stale tags and false-hit the repurposed way; the bypass is now
keyed on {set, filled way}, replacing that way's readout with the
filled line's tag and valid bits.
- MSHR: an allocate could link behind a chain tail dequeuing the same
cycle and be orphaned when its finalize arrived one cycle later;
the dequeuing entry is now excluded from the allocate address match.
Also normalizes parameter-list comments across the cache RTL to
trailing one-liners.
rtlsim: sgemm -n128 8,979,027 -> 8,739,648 cycles, sgemv -5.0%,
sgemm-mc -2.2%; CI cache suite 34/34, amo suite green; VX_cache_top
NT8/NW8 closes 300 MHz on U55C without the AMO engine (the AMO commit
engine remains the pre-existing WNS owner).
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The 2026-07-08 pocl prebuilt refresh switched the SPIR-V translator from the linked LLVMSPIRVLib to an external llvm-spirv, baking the build machine's binary path into libpocl. On any other host the exec fails and every HIP program build aborts with an empty log: chipStar reports hipErrorNotInitialized, the launch no-ops, and all hip CI cells fail with garbage results. OpenCL C cells are unaffected since only SPIR-V IL input goes through llvm-spirv. Set pocl's POCL_PATH_LLVM_SPIRV relocation override to the llvm-vortex bundle's own llvm-spirv, which the run targets' LD_LIBRARY_PATH already covers. With a lib-linked pocl the override is unused, so this is safe across bundle generations. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Fill data no longer rides the pipeline registers: the sector staged at fill accept (the fill buffer) feeds the data-array fill port and the forward-response word directly, so the data-path struct carries a single word (store/AMO payload) instead of a replicated full line, and the write word reaches its commit-stage consumers through a dedicated one-word pipe register. Back-to-back fills are interlocked on deferred data configs (PIPE_EX>0); the classic 2-stage bank needs no interlock. Cycle behavior is bit-identical. NT8/NW8 cache DUT on U55C: 306.7 MHz / 20728 LUT / 27035 FF no-AMO (from 300.8 / 21058 / 28548), 282.8 MHz / 35931 LUT / 32704 FF with AMO (from 277.2 / 35767 / 34279). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A read admitted while an atomic's result is still queued or settling now byte-merges the in-flight writeback bytes over the array word, and a younger plain store to those bytes clears them from the queued entry so the later writeback cannot overwrite it. Closes the window between an atomic's commit and its array write for requests replayed behind it. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The input queue depth doubled as the outstanding-request bound, so independent loads from the same warp serialized on memory instead of overlapping their misses. Track in-flight requests in a dedicated pool (VX_CFG_LSU_PENDING_SIZE, ~2 in-flight loads per warp) while the staging queue keeps its own small depth. The mem-side queue depth is no longer a config knob: its CLOG2 is the dcache-facing tag-id width and the word-coalescer slot count, so it is derived from the pool size in VX_gpu_pkg, with a static assert at the coalescer-bypass boundary where an undersized tag would alias responses. sgemm -31.8%, sgemv -33.7%, sgemm-mc -50.8% cycles; perf gate clean. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Locks in the measured gains (sgemm 0.678x, sgemm-mc 0.493x, sparse-fp16 0.488x, fp16 0.642x, sgemv 0.726x) so the ratchet catches any regression back toward the old numbers. All 27 cases re-measured on rtlsim; instrs unchanged. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Move the cache subsystem, cache flush, and microarchitecture documents into docs/designs and bring them current, adding a full reference for every cache configuration parameter and its performance impact. Rewrite the codebase layout to match the tree, populate references.md with foundational GPU-architecture resources, rename debug_mode.md to kernel_debugging.md, reorganize the documentation index, and update the README (ray-tracing unit, GoogleTest dependency, local CI usage). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Replace the per-(warp,lane) flip-flop slot array with a lane-packed VX_dp_ram (three read mirrors, one fixed-priority write port). Make the RTU bus beat-serial: the ray streams one word per beat straight from the window RAM, and the hit record streams back one word per beat into it, so no ray or hit record is held whole. This removes ray_r, the parallel hit registers, and cb_hit_t_lat; narrows status to a byte; and makes scene_base a per-warp latch plus a scalar bus field. TEX/OM read the window through a runtime slot index on two ports (prologue for the quad-LOD prefetch, one-ahead address for OM), replacing the eight async ports. cons_rd_if narrows 8->2; cons_wr_if gains a ready for texel retire-gating. VX_core_top ties off the graphics cluster buses so dut/core elaborates with RTU/OM/RASTER. Result (dut/core, 8w x 8t, RTU+OM+RASTER, XCU55C, 250 MHz): VX_gfx_window is 638 LUT / 59 FF / 4 BRAM, down from an unsynthesizable flip-flop array; core meets timing (WNS +0.168 ns, Fmax 261 MHz). 30/30 rtlsim raytracing pass; rt_raycast renders bit-exact. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> (cherry picked from commit 1cf47ed99bcae40dcafd2a0b0aaac6ac210b5c1f)
Rewrite xilinx_slr_pblocks.tcl from blind round-robin to a congestion-aware floorplanner with a pre-flight guardrail. - Capacity model uses the platform dynamic-region pblocks (real RM-usable LUTs per SLR), not raw SLR site counts. - Size-aware first-fit-decreasing packing under a per-SLR fill ceiling (SLR_PBLOCKS_FILL_CEILING, default 0.65); anchor pre-loaded with shared caches. - Each extension core is bundled with its private cache so a wide core<->cache bus is never split across a boundary. - Pre-flight guardrail reconciles the atom model against the real netlist and checks per-SLR feasibility; on any failure it applies NO pblocks and lets the automatic SSI partitioner run, instead of a floorplan that dies in Place 30-99. - SLR_PBLOCKS_FORCE=1 to apply despite guardrail (debug). Validated on the graphics 1c post-opt checkpoint (reconcile 99.9%, socket 371K over ceiling -> REJECT -> auto). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Reserve an AMO passthrough-table slot at admission for each probe still in the bank pipe: a probe reaching the pipe head with the table full could not retire, and the full pipe blocked the response processing that frees slots. Defer an atomic while an older same-line access is still inside the bank pipe (its MSHR entry only exists from pipe exit), not just while the MSHR holds one: the atomic could otherwise forward ahead of the older access, letting a younger same-address load chain onto the pre-atomic fill and the fill install a stale line. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
A request admitted during the fill-to-replay handoff probes the MSHR the same cycle a chained atomic dequeues, sees a plain hit, and reads the line before the replayed atomic commits — a same-address ordering violation the commit path cannot catch (the request is ahead of the atomic in the pipe). Drive the LLC branch's req_input_defer from the MSHR's pending-AMO probe, mirroring the non-LLC passthrough guard. The MSHR gains an AMO_PASSTHRU parameter so per-entry AMO tracking now also runs at the LLC while the never-coalesce mask stays passthrough-only, and the probe matches only entries persisted as misses: transient hit-path pre-allocations are already ordered by the commit path, and probing them would serialize hit atomics at the input, starving LR/SC forward progress under contention. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Two functional bugs in the memory-backed window's TEX consumer, exposed by graphics model_parity (gfx_tex / gfx_draw3d) that the RTU-only bring-up missed: 1. The synchronous window read latched each (u,v) pair the cycle its address was driven, but the RAM presents the word one cycle later. `pre_done` fired on the latch cycle, so a single vx_tex4 issued with the PREVIOUS op's coordinates — every textured image shifted by one warp. Gate issue on the settled cycle (pre_pairs + 1). 2. The response-consume handshake (tex_bus_if.rsp_ready) was gated on the window write grant, but the writeback buffer's valid_in was not. While the shared window write port was busy (raster seed / fill contending — never the case in isolated gfx_tex, always in gfx_draw3d), the still-valid response was re-accepted every cycle, writing rd back more than once and tripping the scoreboard's invalid-writeback assertion. Couple valid_in to win_wr_ok. gfx_tex and gfx_draw3d now pass on rtlsim; draw3d retires 143143 instrs, matching the recorded baseline exactly. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
…gister) VX_stream_arb / VX_stream_switch split OUT_BUF internally into the skid (TO_OUT_BUF_SIZE) and the output register (TO_OUT_BUF_REG), so callers must pass the raw OUT_BUF. VX_cache_bypass instead pre-reduced it with TO_OUT_BUF_SIZE, so the value was split twice and the output register was zeroed (TO_OUT_BUF_SIZE(3)=2 then TO_OUT_BUF_REG(2)=0). The cache's outgoing request (mem_bus_out_arb, MEM_OUT_BUF) and response (core_bus_nc_switch, CORE_OUT_BUF) got a skid with no output register. On SSI parts this leaves the L1<->L2 seam combinationally driven across the SLR boundary; with the socket off L2's SLR the path degenerates into a long interposer route and Fmax collapses. Passing raw *_OUT_BUF restores skid+register so the cache hard-registers its external interfaces at the source, both directions. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
All four cases pass execution (correct output on both drivers, retired
instructions match where applicable); the failures are rtlsim<->simx parity
comparisons only, so they xfail rather than fail CI.
- model_parity-tex, model_parity-rt_raycast: cycle-count divergence. The
memory-backed window reads TEX operands synchronously and streams the RTU
bus beat-serially, adding latency the SimX models do not yet reflect. Instrs
match; images are correct/bit-exact. Pending SimX timing alignment.
- model_parity-raster, model_parity-draw3d: pre-existing SimX-vs-RTL
instruction-count divergence. Both drivers render correctly and rtlsim
matches the recorded perf baselines; the SimX raster/graphics fragment model
retires a different instruction count than RTL. Predates this work.
Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
Both L1s forward their flush command down to the shared L2, and a cache that is sweeping locks out incoming core requests for the duration. The icache holds no dirty data, so it retired almost immediately and pushed its flush to the L2 while the dcache was still evicting dirty lines. The L2 then blocked those evictions for its entire sweep, so they reached memory long after the host had already read the buffer back — a write-back L1 over an L2 silently dropped dirty lines (stress dogfood imul, deadbeef results). SimX already drains the hierarchy inside-out (all L1s, then L2, then L3); this brings the RTL in line with the model. Repro (1 core, DCACHE_WRITEBACK=1, --l2cache): Test1-imul lost one 64B line and now passes 22/22. The stress dogfood-1 config (2 clusters, 2 cores, L2+L3, all write-back) passes 22/22 with zero errors. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The TRACE2 head uop allocated its slot from the cluster-shared ray pool inside the SFU. A functional unit retries a uop it cannot complete by leaving it at the head of its input queue, so once the pool was empty that uop jammed the queue — and the WAIT2 stuck behind it was the only thing that could ever release a slot. Deadlock whenever cores x warps exceeds RTU_NUM_CTX, which is why it only ever showed up multi-core (raytracing rt_smoke_mc hung until the job timed out). Claim the slot at issue instead, so the head uop can never stall inside the unit. Claiming rather than merely testing for a free slot also keeps another core from taking the last one between issue and execute. Validated at the default pool size that used to hang: the raytracing category is green (5 passed, 1 xfailed) across both multi-core topologies on simx and rtlsim. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Eight model_parity cases exceed the 5% cycle tolerance (dxa copy and the wgmma/dxa mixes, sgemv, sgemm-mc, sgemmx, tensor fp16). All eight execute correctly and their retired-instruction counts match exactly — the runner asserts instructions before cycles — so the divergence is timing only: SimX does not model the decoupled LSU pending pool (nor, for sgemm-mc, the reworked cache banks), and under-counts against the RTL. This is a documented deviation, not a fix. The real debt is that the RTL LSU landed without the matching SimX timing-model update; modelling that is what actually closes these, and the known_issue markers should come off when it does. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The window's synchronous RAM read costs one cycle per access, moving gfx_draw3d from 454560 to 463873 cycles (+2.0%, just over the gate). Retired instructions are unchanged at 143143, so the workload is the same and only timing moved. Regenerated with --update-baselines. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
toolchain_install.sh never fetches it, nothing under sim/, tests/, hw/, ci/ or sw/ references it, and the test suite is pytest-driven. The only gtest in the tree belongs to the vendored cvfpu flexfloat testbench, which Vortex does not build. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Two width defects made the RTU fail to elaborate at XLEN=64. Both were latent on rtlsim, which builds with -Wno-fatal; the opae and xrt sim builds promote the warnings to errors and so failed outright. scene_base is a device byte address but was fixed at 32 bits, while MEM_ADDR_WIDTH is 48 at XLEN=64. Beyond the width mismatch this truncated the pointer: the gfx window captured it from the low word of rs1. Size it from MEM_ADDR_WIDTH wherever it is declared, packed into a bus payload, or added to an offset, and cast the intra-scene offsets to the address width at the adders. VX_fncp_unit returns an XLEN-wide result -- it also serves the integer-returning compare and class ops -- while its operands are FLEN-wide, and the RTU leaves FLEN at its 32-bit default. The box PE bound that result port directly to fp32 nets. Take the low word through XLEN-wide nets, as the compare result in the same module and the triangle PE already do. At XLEN=32 both widths collapse to 32, so 32-bit builds are unchanged. Co-Authored-By: Claude Opus 4.8 (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
3 participants
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.
Merges the
prismline intomaster. This is a large, long-running branch — 361 commits — so the summary below is organized by area rather than by commit.Hardware
SETW/GETW) used by the RTU, TEX, and OM paths. Any fixed-function stage can be dropped to SIMT software through thegfx_*_swABI, so FF/SW mixes validate against the same golden.Software
vortexpipeGallium driver, includingVK_KHR_ray_query.make installnow lays out a pkg-config sysroot (vortex-runtime.pc/vortex-kernel.pc); downstream tools integrate through$VORTEX_PATHonly.CI 2.0
Replaces the shell-driven regression script with a data-driven pytest harness: test cases live in
ci/testcases/*.yamland are parametrized into pytest items, with markers derived from the data. Addsmodel_parity(SimX vs RTL: instruction counts must match exactly, cycles within tolerance) andperf_gate(cycles against recorded baselines) checks.Notes for review
prismis currently 3 commits behindmaster(b70e0d56a,c9e249a7f,be9ebe05f); those will need to be merged or rebased in before this can land.sim/rtlsimbuilds with-Wno-fatalwhile the opae and xrt sim builds promote Verilator warnings to errors — the latter two are the stricter elaboration gate.🤖 Generated with Claude Code