From 9e96492d8f4467137910e1c46764fe3efbd82be4 Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 02:10:42 -0700 Subject: [PATCH 1/8] cp: decode the AXI-Lite register file flat, not as a priority chain read_reg() and is_decoded() tested every register with its own early return, which synthesizes as a priority chain of ~30 address compares. On ASIC (Yosys + ABC, ASAP7) that chain was the CP's critical path, about 80 logic levels into rresp, and it held the whole `top` DUT to 341 MHz against its 400 MHz target. Both windows decode every word (the 16-byte block padding fills the gaps), so is_decoded() reduces to word alignment plus window membership, and read_reg() becomes one case per window. Yosys proves the new module sequentially equivalent to the old one (equiv_induct, 2445/2445 $equiv cells; the same check flags a one-bit change to the DECERR sentinel). VX_cp_axil_regfile alone 362 -> 1286 MHz, cells 5732 -> 3779 VX_cp_core alone 363 -> 1169 MHz top (8x8, asic_gate) 341 -> 620 MHz Co-Authored-By: Claude Opus 5.5 --- hw/rtl/cp/VX_cp_axil_regfile.sv | 99 +++++++++++++++------------------ 1 file changed, 45 insertions(+), 54 deletions(-) diff --git a/hw/rtl/cp/VX_cp_axil_regfile.sv b/hw/rtl/cp/VX_cp_axil_regfile.sv index 130db04c61..e6fc574b14 100644 --- a/hw/rtl/cp/VX_cp_axil_regfile.sv +++ b/hw/rtl/cp/VX_cp_axil_regfile.sv @@ -182,37 +182,48 @@ module VX_cp_axil_regfile return 1'b1; endfunction + // The global window: 0x000..0x03F, every word of it decoded. + function automatic logic in_global(input logic [ADDR_W-1:0] addr); + return (addr[ADDR_W-1:6] == '0); + endfunction + // ---- Read data combinational decode ---- + // One case per window rather than a chain of early returns: the registers + // are mutually exclusive, and a priority chain of ~30 compares is what + // limited the CP's clock. function automatic logic [31:0] read_reg(input logic [ADDR_W-1:0] addr); logic [QID_W-1:0] qid; logic [5:0] off; - if (is_global(addr, 8'h00)) return r_cp_ctrl; - if (is_global(addr, 8'h04)) return {30'd0, cp_error, cp_busy}; - if (is_global(addr, 8'h08)) return {8'd0, - 8'(AXI_TID_W), - 8'(RING_SIZE_LOG2_MAX), - 8'(NUM_QUEUES)}; - if (is_global(addr, 8'h10)) return r_cycle_count[31:0]; - if (is_global(addr, 8'h14)) return r_cycle_count[63:32]; - if (is_global(addr, 8'h18)) return gpu_dev_caps[31:0]; - if (is_global(addr, 8'h1C)) return gpu_dev_caps[63:32]; - if (is_global(addr, 8'h20)) return gpu_isa_caps[31:0]; - if (is_global(addr, 8'h24)) return gpu_isa_caps[63:32]; - // Block padding -- see the 16-byte block note on is_decoded(). Read as - // zero; they exist so their enclosing block is fully populated. - // 0x28/0x2C are CP_SATP_LO/HI in the software ABI (sw/runtime/common/ - // device.cpp), i.e. the CP DMA MMU page-table root. This CP has no MMU -- - // there is no satp anywhere in hw/rtl/cp -- so it correctly reports - // VM_ENABLED=0 in CP_DEV_CAPS and the runtime never writes them. Reading - // zero is the honest answer; implement them for real if an MMU is added. - if (is_global(addr, 8'h0C)) return 32'd0; - if (is_global(addr, 8'h28)) return 32'd0; - if (is_global(addr, 8'h2C)) return 32'd0; - // AFU host-port drain counters (debug, read-only). - if (is_global(addr, 8'h30)) return dbg_host_w_counts; - if (is_global(addr, 8'h34)) return dbg_host_r_counts; - if (is_global(addr, 8'h38)) return 32'd0; // pad: completes block 0x30 - if (is_global(addr, 8'h3C)) return 32'd0; // pad: completes block 0x30 + if (in_global(addr)) begin + case (addr[5:0]) + 6'h00: return r_cp_ctrl; + 6'h04: return {30'd0, cp_error, cp_busy}; + 6'h08: return {8'd0, + 8'(AXI_TID_W), + 8'(RING_SIZE_LOG2_MAX), + 8'(NUM_QUEUES)}; + 6'h10: return r_cycle_count[31:0]; + 6'h14: return r_cycle_count[63:32]; + 6'h18: return gpu_dev_caps[31:0]; + 6'h1C: return gpu_dev_caps[63:32]; + 6'h20: return gpu_isa_caps[31:0]; + 6'h24: return gpu_isa_caps[63:32]; + // Block padding -- see the 16-byte block note on is_decoded(). Read + // as zero; they exist so their enclosing block is fully populated. + // 0x28/0x2C are CP_SATP_LO/HI in the software ABI (sw/runtime/common/ + // device.cpp), i.e. the CP DMA MMU page-table root. This CP has no + // MMU -- there is no satp anywhere in hw/rtl/cp -- so it correctly + // reports VM_ENABLED=0 in CP_DEV_CAPS and the runtime never writes + // them. Reading zero is the honest answer; implement them for real if + // an MMU is added. + 6'h0C, 6'h28, 6'h2C: return 32'd0; + // AFU host-port drain counters (debug, read-only). + 6'h30: return dbg_host_w_counts; + 6'h34: return dbg_host_r_counts; + 6'h38, 6'h3C: return 32'd0; // pad: completes block 0x30 + default: return 32'hDEAD_BEEF; // unaligned: DECERR + endcase + end if (decode_queue(addr, qid, off)) begin case (off) 6'h00: return r_ring_base[qid][31:0]; @@ -255,37 +266,17 @@ module VX_cp_axil_regfile // block worked -- it was never about the registers, only their blocks. // // Consequence: when adding a register, pad its block out to 4 words. + // + // With the padding, every word of the global window and of each queue's + // 0x40 block is decoded (the queue's 0x34..0x3C pad completes its 0x30 + // block so Q_LAST_DCR_RSP is reachable), so the test is alignment plus + // window membership -- flat logic, not a compare per register. function automatic logic is_decoded(input logic [ADDR_W-1:0] addr); logic [QID_W-1:0] qid; // populated by decode_queue but unused here logic [5:0] off; `UNUSED_VAR (qid) - if (is_global(addr, 8'h00)) return 1'b1; - if (is_global(addr, 8'h04)) return 1'b1; - if (is_global(addr, 8'h08)) return 1'b1; - if (is_global(addr, 8'h0C)) return 1'b1; // pad: completes block 0x00 - if (is_global(addr, 8'h10)) return 1'b1; - if (is_global(addr, 8'h14)) return 1'b1; - if (is_global(addr, 8'h18)) return 1'b1; - if (is_global(addr, 8'h1C)) return 1'b1; - if (is_global(addr, 8'h20)) return 1'b1; - if (is_global(addr, 8'h24)) return 1'b1; - if (is_global(addr, 8'h28)) return 1'b1; // pad: completes block 0x20 - if (is_global(addr, 8'h2C)) return 1'b1; // pad: completes block 0x20 - if (is_global(addr, 8'h30)) return 1'b1; // drain counters (debug) - if (is_global(addr, 8'h34)) return 1'b1; - if (is_global(addr, 8'h38)) return 1'b1; // pad: completes block 0x30 - if (is_global(addr, 8'h3C)) return 1'b1; // pad: completes block 0x30 - if (decode_queue(addr, qid, off)) begin - case (off) - 6'h00, 6'h04, 6'h08, 6'h0C, 6'h10, 6'h14, - 6'h18, 6'h1C, 6'h20, 6'h24, 6'h28, 6'h2C, - 6'h30, - // pad: completes the queue's 0x30 block so Q_LAST_DCR_RSP is reachable - 6'h34, 6'h38, 6'h3C: return 1'b1; - default: return 1'b0; - endcase - end - return 1'b0; + `UNUSED_VAR (off) + return (addr[1:0] == 2'b00) && (in_global(addr) || decode_queue(addr, qid, off)); endfunction // ============================================================================ From 7518c88783fcb68ea9a2ca92d0afcdbfab85d63c Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 02:10:43 -0700 Subject: [PATCH 2/8] raster: carry the early-Z depth plane through the quad-evaluator stage The block evaluator fed its output FIFO zplane_r from the stage before VX_raster_qe, while the stamps come from VX_raster_qe's registered output. Each wave therefore left with the depth plane of the block behind it, and whenever consecutive blocks belonged to different primitives, early-Z evaluated a quad against the wrong triangle's plane. With the wrong candidate depth it culled a fragment that late-Z keeps, so the image differed: gfx_draw3d-6 (evilskull 128) failed with early-Z on rtlsim, 2 pixels at 2 slices and 1 pixel at 1 slice, while SimX passed. The plane now goes through a register with the evaluator's enable, so each wave carries its own primitive's plane. Early-Z on, rtlsim: gfx_draw3d-6 passes at 1 and 2 slices, matching SimX. The early-Z cases in graphics.yaml never caught this: they passed -DVX_CFG_RASTER_EARLYZ, which is not a config knob, so they built the knob-off core. Co-Authored-By: Claude Opus 5.5 --- hw/rtl/raster/VX_raster_be.sv | 19 ++++++++++++++++++- 1 file changed, 18 insertions(+), 1 deletion(-) diff --git a/hw/rtl/raster/VX_raster_be.sv b/hw/rtl/raster/VX_raster_be.sv index 7c68132f63..e33ca50d9c 100644 --- a/hw/rtl/raster/VX_raster_be.sv +++ b/hw/rtl/raster/VX_raster_be.sv @@ -147,6 +147,23 @@ module VX_raster_be import VX_raster_pkg::*; #( .yloc_out (qe_yloc) ); +`ifdef VX_CFG_RASTER_EARLYZ_ENABLE + // The quad evaluator registers its outputs once; delay the plane the same + // way so each wave leaves with its own primitive's plane, not the plane of + // the block behind it. + wire [2:0][`RASTER_DATA_BITS-1:0] qe_zplane; + + VX_pipe_register #( + .DATAW (3 * `RASTER_DATA_BITS) + ) qe_zplane_reg ( + .clk (clk), + .reset (reset), + .enable (~stall), + .data_in (zplane_r), + .data_out (qe_zplane) + ); +`endif + // Populate fifo inputs wire [OUTPUT_BATCHES-1:0][OUTPUT_QUADS-1:0] fifo_mask_in; @@ -227,7 +244,7 @@ module VX_raster_be import VX_raster_pkg::*; #( .reset (reset), .valid_in (fifo_valid_in), .ready_in (fifo_ready_in), - .data_in ({zplane_r, fifo_stamp_in[fifo_arb_index]}), + .data_in ({qe_zplane, fifo_stamp_in[fifo_arb_index]}), .data_out ({zplane_out, stamps_out}), .valid_out (valid_out), .ready_out (ready_out) From fec222407a6e2f4b69acd22b0553c45717f1293c Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 02:10:54 -0700 Subject: [PATCH 3/8] config: enable raster early-Z wherever RASTER and OM are both present VX_CFG_RASTER_EARLYZ_ENABLE now defaults to (EXT_RASTER && EXT_OM) instead of false. Early-Z reads committed depth through the OM's ocache, so it is legal only with both units (VX_define.vh errors otherwise); the expression keeps every other build unchanged, and -DVX_CFG_RASTER_EARLYZ_DISABLE opts a build out. The early-Z cases in graphics.yaml passed -DVX_CFG_RASTER_EARLYZ, which resolves to nothing, so CI never ran the early-Z path. Every RASTER+OM case now covers it; those five cases become gfx_no_earlyz-* and pass -DVX_CFG_RASTER_EARLYZ_DISABLE to keep the knob-off build covered. Validated with early-Z on (after the be.sv plane fix): graphics simx 35/35, graphics_parity simx 6/6, vulkan simx 13/13, graphics rtlsim 43/43, graphics_parity rtlsim 6/6, vulkan xrtsim 1/1. The generated config header changed, so every perf baseline's config_hash is re-recorded. Cycles are identical outside graphics; draw3d moves by -1.61% (nt16, +1041 instrs), +0.03% (nt4) and +0.02% (mc). Co-Authored-By: Claude Opus 5.5 --- VX_config.toml | 4 +++- ci/baselines/perf/core.json | 32 ++++++++++++------------- ci/baselines/perf/dxa.json | 22 ++++++++--------- ci/baselines/perf/graphics.json | 26 ++++++++++----------- ci/baselines/perf/raytracing.json | 6 ++--- ci/baselines/perf/tensor.json | 10 ++++---- ci/baselines/perf/tensor_mx.json | 2 +- ci/baselines/perf/tensor_sp.json | 2 +- ci/baselines/perf/tensor_wg.json | 6 ++--- ci/testcases/graphics.yaml | 39 ++++++++++++++++--------------- 10 files changed, 76 insertions(+), 73 deletions(-) diff --git a/VX_config.toml b/VX_config.toml index 71b9c59b73..a3336551a8 100644 --- a/VX_config.toml +++ b/VX_config.toml @@ -375,7 +375,9 @@ VX_CFG_RASTER_BLOCK_LOG_SIZE = 2 VX_CFG_RASTER_MEM_FIFO_DEPTH = 8 VX_CFG_RASTER_QUAD_FIFO_DEPTH = 32 VX_CFG_RASTER_MEM_QUEUE_SIZE = 4 -VX_CFG_RASTER_EARLYZ_ENABLE = false +# Early-Z occlusion cull: on wherever it is legal. It reads committed depth +# through the OM's ocache, so it needs both units (VX_define.vh enforces it). +VX_CFG_RASTER_EARLYZ_ENABLE = "expr: $VX_CFG_EXT_RASTER_ENABLE and $VX_CFG_EXT_OM_ENABLE" [om] VX_CFG_NUM_OM_CORES = "expr: max(1, up($VX_CFG_NUM_CORES / 8))" diff --git a/ci/baselines/perf/core.json b/ci/baselines/perf/core.json index fb6ef0491c..5e226eba16 100644 --- a/ci/baselines/perf/core.json +++ b/ci/baselines/perf/core.json @@ -6,7 +6,7 @@ }, "app": "raycast", "args": "", - "config_hash": "ae710a554c1b4afe", + "config_hash": "ee3a227211fa9942", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "raycast", "args": "", - "config_hash": "204dd696e55b22c0", + "config_hash": "c8b680eabeb18655", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -28,7 +28,7 @@ }, "app": "sgemm", "args": "-n128", - "config_hash": "4e0d1121c7c753fb", + "config_hash": "58eb3cc06964207a", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=16 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, @@ -39,7 +39,7 @@ }, "app": "sgemm", "args": "-n128", - "config_hash": "2e368e29b49794c1", + "config_hash": "aa6f26c7306417a8", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, @@ -50,7 +50,7 @@ }, "app": "sgemm", "args": "-n128", - "config_hash": "8493bfcbea192f2b", + "config_hash": "b390d957eec8220f", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -61,7 +61,7 @@ }, "app": "sgemm", "args": "-n128", - "config_hash": "ef1d8528f632523b", + "config_hash": "13d03df76d55105e", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -72,7 +72,7 @@ }, "app": "sgemmx", "args": "-n128", - "config_hash": "c016962fb7669255", + "config_hash": "68638b6685801699", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -83,7 +83,7 @@ }, "app": "sgemmx", "args": "-n128", - "config_hash": "65047df50cd355e6", + "config_hash": "837240be3049d62e", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -94,7 +94,7 @@ }, "app": "sgemv", "args": "-m512 -n512", - "config_hash": "70d3eb40917f6df3", + "config_hash": "c6a3720118661bf7", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -105,7 +105,7 @@ }, "app": "sgemv", "args": "-m512 -n512", - "config_hash": "e359f1a79bc49193", + "config_hash": "db6bb5af90ecc1a5", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -116,7 +116,7 @@ }, "app": "softmax", "args": "-n64", - "config_hash": "91488cd182a7671a", + "config_hash": "1c0e7742deb14093", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -127,7 +127,7 @@ }, "app": "softmax", "args": "-n64", - "config_hash": "c8061d3015fc6d7e", + "config_hash": "3797eecae3200315", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -138,7 +138,7 @@ }, "app": "stencil3d", "args": "-n32", - "config_hash": "e659f36c50fcf6bf", + "config_hash": "aa768c2a9c65b092", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -149,7 +149,7 @@ }, "app": "stencil3d", "args": "-n32", - "config_hash": "3bbd7b84bf73c746", + "config_hash": "5400d882a6c96ac1", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -160,7 +160,7 @@ }, "app": "vecadd", "args": "-n16384", - "config_hash": "a5ce692bba42344b", + "config_hash": "98b8bdc5e546ea62", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -171,7 +171,7 @@ }, "app": "vecadd", "args": "-n16384", - "config_hash": "7ad7759234a8fedf", + "config_hash": "1510b8438164b252", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" } diff --git a/ci/baselines/perf/dxa.json b/ci/baselines/perf/dxa.json index 64f8cb0c2b..d864f76b96 100644 --- a/ci/baselines/perf/dxa.json +++ b/ci/baselines/perf/dxa.json @@ -6,7 +6,7 @@ }, "app": "dxa_copy_mcast", "args": "", - "config_hash": "d8dc71acb72624de", + "config_hash": "cbcd10be884ceab7", "configs": "-DVX_CFG_EXT_DXA_ENABLE", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "dxa_copy", "args": "-d1", - "config_hash": "84239768cfc34b0d", + "config_hash": "3a1ec1b81b5b3baf", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -28,7 +28,7 @@ }, "app": "dxa_copy", "args": "-d1", - "config_hash": "e315513bdc6942fb", + "config_hash": "3a888067fcc449d6", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -39,7 +39,7 @@ }, "app": "sgemm2_dxa_mcast", "args": "-n128", - "config_hash": "d6edd97b6570efdd", + "config_hash": "14050e4dcef65266", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -50,7 +50,7 @@ }, "app": "sgemm2_dxa_mcast", "args": "-n128", - "config_hash": "f638e7687fcfdf72", + "config_hash": "a55b900d67aa3657", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -61,7 +61,7 @@ }, "app": "sgemm2_dxa", "args": "-n128 -t4 -m1", - "config_hash": "538516a52b3616f9", + "config_hash": "36f5886fdc97e040", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -72,7 +72,7 @@ }, "app": "sgemm2_dxa", "args": "-n128 -t4 -m1", - "config_hash": "b68ce24c030ee769", + "config_hash": "9ccf890d5ac45a62", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -83,7 +83,7 @@ }, "app": "sgemm_tcu_wg_dxa_mcast", "args": "-m 128 -n 128 -k 64", - "config_hash": "f2390772053e2d69", + "config_hash": "ee7678a7e7e9b79b", "configs": "-DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1 -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -94,7 +94,7 @@ }, "app": "sgemm_tcu_wg_dxa_mcast", "args": "-m 128 -n 128 -k 64", - "config_hash": "66b8301aa43d2744", + "config_hash": "c79015e4caec7b93", "configs": "-DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1 -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -105,7 +105,7 @@ }, "app": "sgemm_tcu_wg_dxa", "args": "-m 128 -n 128 -k 64", - "config_hash": "ba8978d35dc6e582", + "config_hash": "5a35710ef3aec3a6", "configs": "-DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DWGMMA_NRC=16 -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_THREADS=32", "driver": "rtlsim" }, @@ -116,7 +116,7 @@ }, "app": "sgemm_tcu_wg_sp_dxa", "args": "-m 128 -n 128 -k 128", - "config_hash": "9480be10a4650239", + "config_hash": "d12fad3fd43e61c6", "configs": "-DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_TCU_SPARSE_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DWGMMA_NRC=32 -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_THREADS=32", "driver": "rtlsim" } diff --git a/ci/baselines/perf/graphics.json b/ci/baselines/perf/graphics.json index e67ff8d682..32c723239d 100644 --- a/ci/baselines/perf/graphics.json +++ b/ci/baselines/perf/graphics.json @@ -1,34 +1,34 @@ { "graphics:perf_gate-draw3d-mc:rtlsim": { "32": { - "cycles": 531658, + "cycles": 531740, "instrs": 204102 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "2647c06a6e2826bb", + "config_hash": "c8c2af44277e2cb0", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, "graphics:perf_gate-draw3d-nt16:rtlsim": { "32": { - "cycles": 148046, - "instrs": 51475 + "cycles": 145662, + "instrs": 52516 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "a4c096131a6399a9", + "config_hash": "ee0410ff09e7ea58", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "graphics:perf_gate-draw3d-nt4:rtlsim": { "32": { - "cycles": 531210, + "cycles": 531359, "instrs": 204102 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "ef55e8acfb29082e", + "config_hash": "00c7939a6d958e0e", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -39,7 +39,7 @@ }, "app": "gfx_om", "args": "-w 128 -h 128 -r whitebox_128.png", - "config_hash": "6f583da43550b85f", + "config_hash": "60103e7b6eb2c625", "configs": "-DVX_CFG_EXT_OM_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -50,7 +50,7 @@ }, "app": "gfx_om", "args": "-w 128 -h 128 -r whitebox_128.png", - "config_hash": "adf8fad7a766960a", + "config_hash": "ee88907a281d364a", "configs": "-DVX_CFG_EXT_OM_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -61,7 +61,7 @@ }, "app": "gfx_raster", "args": "-w 128 -h 128 -t triangle.cgltrace -r triangle_ref_128.png", - "config_hash": "b989d0754c6c0eb4", + "config_hash": "3ba867af954e5cbf", "configs": "-DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -72,7 +72,7 @@ }, "app": "gfx_raster", "args": "-w 128 -h 128 -t triangle.cgltrace -r triangle_ref_128.png", - "config_hash": "1d94de7ac5d4207f", + "config_hash": "859b3c4bd2208b92", "configs": "-DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -83,7 +83,7 @@ }, "app": "gfx_tex", "args": "-i toad.png -r toad_ref_f0.png -f 0 -g 0", - "config_hash": "4d4b6feec7071948", + "config_hash": "1c38a250b6b85273", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -94,7 +94,7 @@ }, "app": "gfx_tex", "args": "-i toad.png -r toad_ref_f0.png -f 0 -g 0", - "config_hash": "9a76bd5faf821bd1", + "config_hash": "226df03058dfb436", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" } diff --git a/ci/baselines/perf/raytracing.json b/ci/baselines/perf/raytracing.json index c276c9582e..86b562e3b1 100644 --- a/ci/baselines/perf/raytracing.json +++ b/ci/baselines/perf/raytracing.json @@ -6,7 +6,7 @@ }, "app": "rt_raycast", "args": "", - "config_hash": "e4a96d9dbbed20d6", + "config_hash": "c71b720ea7aa7e72", "configs": "-DVX_CFG_EXT_RTU_ENABLE -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "rt_raycast", "args": "", - "config_hash": "39e834ab9719b706", + "config_hash": "3049a9d721044cde", "configs": "-DVX_CFG_EXT_RTU_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -28,7 +28,7 @@ }, "app": "rt_raycast", "args": "", - "config_hash": "268349756ebd862c", + "config_hash": "14ae0ed831d1cd4a", "configs": "-DVX_CFG_EXT_RTU_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" } diff --git a/ci/baselines/perf/tensor.json b/ci/baselines/perf/tensor.json index 7e88c1cc60..3249cc1b16 100644 --- a/ci/baselines/perf/tensor.json +++ b/ci/baselines/perf/tensor.json @@ -6,7 +6,7 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "057e7fd450c175d3", + "config_hash": "6a095ed165469a30", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32 -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "a2e90c72f59e67af", + "config_hash": "0fe2a8e2a75fcc10", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=16 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, @@ -28,7 +28,7 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "1cac233a0ec6e843", + "config_hash": "b2ef81f2fa1e6a43", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, @@ -39,7 +39,7 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "8f690df8411c6c17", + "config_hash": "5534ddd0051a5c15", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32", "driver": "rtlsim" }, @@ -50,7 +50,7 @@ }, "app": "sgemm2_tcu", "args": "-m 128 -n 128 -k 64", - "config_hash": "0a1788de422861bd", + "config_hash": "421a4ab52f0f072a", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32 -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1", "driver": "rtlsim" } diff --git a/ci/baselines/perf/tensor_mx.json b/ci/baselines/perf/tensor_mx.json index f5545bc8b5..c9b137227e 100644 --- a/ci/baselines/perf/tensor_mx.json +++ b/ci/baselines/perf/tensor_mx.json @@ -6,7 +6,7 @@ }, "app": "sgemm_tcu_mx", "args": "-m 128 -n 128 -k32", - "config_hash": "83cc59b6cf35e0ee", + "config_hash": "719fe3a7ce81115d", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_MX_ENABLE -DVX_CFG_TCU_FP8_ENABLE -DITYPE=mxfp8 -DOTYPE=fp32 -DVX_CFG_TCU_TYPE_DPI", "driver": "rtlsim" } diff --git a/ci/baselines/perf/tensor_sp.json b/ci/baselines/perf/tensor_sp.json index f14e6f8384..e8e15a7ab2 100644 --- a/ci/baselines/perf/tensor_sp.json +++ b/ci/baselines/perf/tensor_sp.json @@ -6,7 +6,7 @@ }, "app": "sgemm_tcu_sp", "args": "-m 128 -n 128 -k 128", - "config_hash": "4f2460323b85e99a", + "config_hash": "acf4fa4675243f0f", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_SPARSE_ENABLE", "driver": "rtlsim" } diff --git a/ci/baselines/perf/tensor_wg.json b/ci/baselines/perf/tensor_wg.json index c8a6b5f353..6ea2c1d004 100644 --- a/ci/baselines/perf/tensor_wg.json +++ b/ci/baselines/perf/tensor_wg.json @@ -6,7 +6,7 @@ }, "app": "sgemm_tcu_wg", "args": "-m64 -n64 -k64", - "config_hash": "5a0e00692a609e17", + "config_hash": "7e8b9759dc36dba0", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1 -DVX_CFG_TCU_FEDP2K -DITYPE=fp16 -DOTYPE=fp32 -DWGMMA_NRC=8 -DWGMMA_RS", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "sgemm_tcu_wg", "args": "-m 128 -n 128 -k 128", - "config_hash": "626ecbcf48b8d1fa", + "config_hash": "c5a71c297fc4dafe", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DITYPE=fp16 -DOTYPE=fp32 -DWGMMA_NRC=32 -DWGMMA_SS -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1", "driver": "rtlsim" }, @@ -28,7 +28,7 @@ }, "app": "sgemm_tcu_wg_sp", "args": "-m 128 -n 128 -k 128", - "config_hash": "bbdf6028e5c7190a", + "config_hash": "ab1451598dfbc872", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8 -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_TCU_SPARSE_ENABLE -DWGMMA_NRC=32 -DVX_CFG_NUM_ALU_BLOCKS=1 -DVX_CFG_NUM_LSU_BLOCKS=1 -DVX_CFG_NUM_FPU_BLOCKS=1", "driver": "rtlsim" } diff --git a/ci/testcases/graphics.yaml b/ci/testcases/graphics.yaml index 7214c48ba0..704debca65 100644 --- a/ci/testcases/graphics.yaml +++ b/ci/testcases/graphics.yaml @@ -472,50 +472,51 @@ tests: threads: 4 args: -w 32 -h 32 -t evilskull.cgltrace -r evilskull_ref_32.png -# ── Early-Z (occlusion cull) — VX_CFG_RASTER_EARLYZ knob ON ─────────────────── -# The raster unit reads committed depth (coherent ocache) and culls occluded -# fragments before shading; requires the OM + RASTER fixed-function units. The -# fixed-point depth-plane contract makes the cull bit-identical to the -# no-early-Z golden (the per-draw VX_DCR_OM_EARLYZ_SAFE runtime gate is set by -# the driver). Knob-off leaves the raster core byte-identical (covered above). -- id: gfx_earlyz-box-simx +# ── Early-Z (occlusion cull) knob OFF ───────────────────────────────────────── +# Early-Z is on by default wherever RASTER + OM are both enabled, so every +# RASTER+OM case above runs with it: the raster unit reads committed depth +# (coherent ocache) and culls occluded fragments before shading, gated per draw +# by VX_DCR_OM_EARLYZ_SAFE, which the driver sets. The fixed-point depth-plane +# contract makes the cull bit-identical to the ROP-only path, so both share one +# golden. These cases keep the knob-off build covered. +- id: gfx_no_earlyz-box-simx via: blackbox drivers: [simx] app: gfx_draw3d - configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ + configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ_DISABLE args: -w 128 -h 128 -t box.cgltrace -r box_ref_128.png touches: [hw/rtl/raster/, sim/simx/raster/, sw/common/gfx_setup.h, sw/kernel/include/vx_graphics.h] -- id: gfx_earlyz-box-rtlsim +- id: gfx_no_earlyz-box-rtlsim via: blackbox drivers: [rtlsim] app: gfx_draw3d - configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ + configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ_DISABLE args: -w 128 -h 128 -t box.cgltrace -r box_ref_128.png touches: [hw/rtl/raster/, sim/simx/raster/, sw/common/gfx_setup.h, sw/kernel/include/vx_graphics.h] -- id: gfx_earlyz-evilskull32-simx +- id: gfx_no_earlyz-evilskull32-simx via: blackbox drivers: [simx] app: gfx_draw3d - configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ -DICACHE_DISABLE -DDCACHE_DISABLE -DLMEM_DISABLE + configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ_DISABLE -DICACHE_DISABLE -DDCACHE_DISABLE -DLMEM_DISABLE shape: {threads: 4} args: -w 32 -h 32 -t evilskull.cgltrace -r evilskull_ref_32.png touches: [hw/rtl/raster/, sim/simx/raster/, sw/common/gfx_setup.h, sw/kernel/include/vx_graphics.h] -- id: gfx_earlyz-evilskull32-rtlsim +- id: gfx_no_earlyz-evilskull32-rtlsim via: blackbox drivers: [rtlsim] app: gfx_draw3d - configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ -DICACHE_DISABLE -DDCACHE_DISABLE -DLMEM_DISABLE + configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ_DISABLE -DICACHE_DISABLE -DDCACHE_DISABLE -DLMEM_DISABLE shape: {threads: 4} args: -w 32 -h 32 -t evilskull.cgltrace -r evilskull_ref_32.png touches: [hw/rtl/raster/, sim/simx/raster/, sw/common/gfx_setup.h, sw/kernel/include/vx_graphics.h] -# evilskull@128 with early-Z on rtlsim: the early-Z cull uses the reflexive -# relaxation of the depth func (drop only strictly-behind fragments), so it is -# image-identical to the ROP-only path (the OM late-Z stays authoritative). -- id: gfx_earlyz-evilskull128-rtlsim +# evilskull@128 without early-Z on rtlsim: the early-Z cull uses the reflexive +# relaxation of the depth func (drop only strictly-behind fragments), so the +# ROP-only path must match the same golden (the OM late-Z stays authoritative). +- id: gfx_no_earlyz-evilskull128-rtlsim via: blackbox drivers: [rtlsim] app: gfx_draw3d - configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ + configs: -DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_EXT_OM_ENABLE -DVX_CFG_RASTER_EARLYZ_DISABLE args: -w 128 -h 128 -t evilskull.cgltrace -r evilskull_ref_128.png tier: full touches: [hw/rtl/raster/, sw/common/gfx_setup.h] From 27312cbfd8d3dd6f52b75593230c9c31a52fb79d Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 11:13:41 -0700 Subject: [PATCH 4/8] ci: gate asic top and gfx at NT=NW=4 and record the four new baselines top, gfx, tensor and vm were added to asic_gate in 4966e76eb at NT=NW=8 without baselines, so the nightly reported them STALE -- and lost top and gfx outright: at 8x8 their Yosys peaks (18.6 and 19.5 GB) exceed a hosted runner's 16 GB, and run 37263315588 killed both mid-synthesis (exit 143). top and gfx are gated at NT=NW=4, where they fit; tensor and vm stay at 8x8. Recorded with the CI toolchain (Yosys 0.69, OpenSTA 3.1.0, sv2v 0.0.13-13), 400 MHz target, on this branch's RTL (top includes the CP regfile fix, gfx the early-Z plane fix): top 4x4 620.9 MHz 1.83M cells 11.2 GB peak RSS gfx 4x4 850.0 MHz 2.31M cells 14.6 GB peak RSS tensor 8x8 976.9 MHz 1.01M cells 8.0 GB peak RSS vm 8x8 1028.5 MHz 1.37M cells 5.3 GB peak RSS Co-Authored-By: Claude Opus 5.5 --- ci/baselines/synthesis/yosys/core.json | 26 ++++++++++++++++++++++ ci/baselines/synthesis/yosys/graphics.json | 26 +++++++++++----------- ci/baselines/synthesis/yosys/tensor.json | 26 +++++++++++----------- ci/baselines/synthesis/yosys/vm.json | 26 +++++++++++----------- ci/testcases/asic_gate.yaml | 19 +++++++++------- 5 files changed, 76 insertions(+), 47 deletions(-) diff --git a/ci/baselines/synthesis/yosys/core.json b/ci/baselines/synthesis/yosys/core.json index 11a2f6a560..9449d60c6b 100644 --- a/ci/baselines/synthesis/yosys/core.json +++ b/ci/baselines/synthesis/yosys/core.json @@ -50,5 +50,31 @@ "tns_ns": 0.0, "wns_ns": 0.796 } + }, + "top": { + "clock_mhz": 400, + "config_hash": "0aa70ccdf8cec0ef", + "configs": "-DVX_CFG_NUM_THREADS=4 -DVX_CFG_NUM_WARPS=4", + "dut": "top", + "result": { + "build_time_s": 2192, + "cell_area_um2": 187229.507, + "cell_count": 1830007, + "env": { + "corner": "tt", + "pdk": "asap7", + "sta": "3.1.0", + "sv2v": "sv2v v0.0.13-13-g493a88f", + "vt": "rvt", + "xlen": 32, + "yosys": "Yosys 0.69" + }, + "fmax_mhz": 620.887, + "power_mw": 126.0, + "seq_area_um2": 71938.595, + "sram_area_um2": 94420.8, + "tns_ns": 0.0, + "wns_ns": 0.889 + } } } diff --git a/ci/baselines/synthesis/yosys/graphics.json b/ci/baselines/synthesis/yosys/graphics.json index 9b7950faa6..e68752040c 100644 --- a/ci/baselines/synthesis/yosys/graphics.json +++ b/ci/baselines/synthesis/yosys/graphics.json @@ -1,28 +1,28 @@ { "gfx": { "clock_mhz": 400, - "config_hash": "719c27411ebd01c9", - "configs": "", + "config_hash": "a5aad82cf4a5712f", + "configs": "-DVX_CFG_NUM_THREADS=4 -DVX_CFG_NUM_WARPS=4", "dut": "gfx", "result": { - "build_time_s": 3920, - "cell_area_um2": 395087.746, - "cell_count": 3785747, + "build_time_s": 2242, + "cell_area_um2": 242202.829, + "cell_count": 2305942, "env": { "corner": "tt", "pdk": "asap7", - "sta": "2.7.0", - "sv2v": "sv2v v0.0.13-3-g80a2f0c", + "sta": "3.1.0", + "sv2v": "sv2v v0.0.13-13-g493a88f", "vt": "rvt", "xlen": 32, - "yosys": "Yosys 0.57+157" + "yosys": "Yosys 0.69" }, - "fmax_mhz": 842.105, - "power_mw": 1100.0, - "seq_area_um2": 151023.722, - "sram_area_um2": 322876.0, + "fmax_mhz": 849.979, + "power_mw": 624.0, + "seq_area_um2": 95716.534, + "sram_area_um2": 307526.4, "tns_ns": 0.0, - "wns_ns": 1.312 + "wns_ns": 1.323 } }, "om": { diff --git a/ci/baselines/synthesis/yosys/tensor.json b/ci/baselines/synthesis/yosys/tensor.json index 388bc1c9c9..9c1ace23cb 100644 --- a/ci/baselines/synthesis/yosys/tensor.json +++ b/ci/baselines/synthesis/yosys/tensor.json @@ -27,28 +27,28 @@ }, "tensor": { "clock_mhz": 400, - "config_hash": "d28e96697d5810ba", - "configs": "", + "config_hash": "7a067e4ffbc02de1", + "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8", "dut": "tensor", "result": { - "build_time_s": 2880, - "cell_area_um2": 224202.142, - "cell_count": 2351511, + "build_time_s": 1086, + "cell_area_um2": 104438.042, + "cell_count": 1014527, "env": { "corner": "tt", "pdk": "asap7", - "sta": "2.7.0", - "sv2v": "sv2v v0.0.13-3-g80a2f0c", + "sta": "3.1.0", + "sv2v": "sv2v v0.0.13-13-g493a88f", "vt": "rvt", "xlen": 32, - "yosys": "Yosys 0.57+157" + "yosys": "Yosys 0.69" }, - "fmax_mhz": 926.012, - "power_mw": 237.0, - "seq_area_um2": 70778.318, - "sram_area_um2": 257618.4, + "fmax_mhz": 976.944, + "power_mw": 122.0, + "seq_area_um2": 40998.085, + "sram_area_um2": 255218.4, "tns_ns": 0.0, - "wns_ns": 1.42 + "wns_ns": 1.476 } } } diff --git a/ci/baselines/synthesis/yosys/vm.json b/ci/baselines/synthesis/yosys/vm.json index 554d178116..07351dcbdd 100644 --- a/ci/baselines/synthesis/yosys/vm.json +++ b/ci/baselines/synthesis/yosys/vm.json @@ -1,28 +1,28 @@ { "vm": { "clock_mhz": 400, - "config_hash": "87c316a577a47ae1", - "configs": "", + "config_hash": "6437095ffd66d32f", + "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8", "dut": "vm", "result": { - "build_time_s": 680, - "cell_area_um2": 103431.322, - "cell_count": 1007116, + "build_time_s": 946, + "cell_area_um2": 146445.005, + "cell_count": 1368872, "env": { "corner": "tt", "pdk": "asap7", - "sta": "2.7.0", - "sv2v": "sv2v v0.0.13-3-g80a2f0c", + "sta": "3.1.0", + "sv2v": "sv2v v0.0.13-13-g493a88f", "vt": "rvt", "xlen": 32, - "yosys": "Yosys 0.57+157" + "yosys": "Yosys 0.69" }, - "fmax_mhz": 973.425, - "power_mw": 134.0, - "seq_area_um2": 43729.211, - "sram_area_um2": 60052.8, + "fmax_mhz": 1028.489, + "power_mw": 185.0, + "seq_area_um2": 67394.009, + "sram_area_um2": 272849.6, "tns_ns": 0.0, - "wns_ns": 1.473 + "wns_ns": 1.528 } } } diff --git a/ci/testcases/asic_gate.yaml b/ci/testcases/asic_gate.yaml index 49a1f5d63b..8a8c0dd34a 100644 --- a/ci/testcases/asic_gate.yaml +++ b/ci/testcases/asic_gate.yaml @@ -33,14 +33,15 @@ # and a divergence between them is a real finding, not an artefact of two lists. # # THIS LIST IS BOUNDED BY MEMORY. It gates the same DUTs as fpga_gate.yaml, but -# the wide ones at NT=NW=8 instead of 16. A hosted runner has 16 GB of RAM; -# Yosys holds the whole elaborated netlist and offers no cap to tune, so a DUT -# that needs more does not fail -- the VM is killed under it and the job ends at +# the wide ones at a reduced NT=NW. A hosted runner has 16 GB of RAM; Yosys +# holds the whole elaborated netlist and offers no cap to tune, so a DUT that +# needs more does not fail -- the VM is killed under it and the job ends at # exit 143 with no report, which looks exactly like an infrastructure outage. At # 16/16, `gfx` measured a 44.2 GB peak RSS for a single Yosys process (3.76M -# cells, 100 min wall), and `core` and `tex` measured 20.7 GB and 28.0 GB. -# Halving NT and NW quarters the lane x warp product that carries that cost. -# fpga_gate.yaml still gates the full-width configurations. +# cells, 100 min wall), and `core` and `tex` measured 20.7 GB and 28.0 GB; +# NT and NW size the lane x warp product that carries that cost, so the wide +# DUTs are gated at NT=NW=8, and `top` and `gfx` at NT=NW=4: at 8/8 they peak +# at 18.6 and 19.5 GB. fpga_gate.yaml gates the full-width configurations. # # `vm` once passed a hosted nightly at 1.0M cells in 11 minutes only because its # L2 was a BYPASS: VX_vm_top wires the L2 as @@ -122,7 +123,8 @@ builds: group: core dut: top clock_mhz: 400 - configs: "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8" + # 11.2 GB Yosys peak at NT=NW=4 (8/8: 18.6 GB). + configs: "-DVX_CFG_NUM_THREADS=4 -DVX_CFG_NUM_WARPS=4" # --- dxa --------------------------------------------------------------------- @@ -188,7 +190,8 @@ builds: group: graphics dut: gfx clock_mhz: 400 - configs: "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_NUM_WARPS=8" + # 14.6 GB Yosys peak at NT=NW=4 (8/8: 19.5 GB). + configs: "-DVX_CFG_NUM_THREADS=4 -DVX_CFG_NUM_WARPS=4" # --- vm ---------------------------------------------------------------------- From 48c972692484498a3c4e127e431552fa4038e31d Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 02:11:04 -0700 Subject: [PATCH 5/8] ci: shard long cells on recorded runtimes, 3 h timeout, ccache back on Timeout and sharding. The tests job times out at 180 minutes (was 300). A cell whose recorded runtime exceeds SHARD_BUDGET_S (150 min) is planned as several shards, split longest-first onto the least-loaded shard by the per-case wall times in ci/runtimes.json; `pytest --shard=K/N` derives the same split from the same ids. ci/record_runtimes.py refreshes the table from a CI run's junit artifacts (gh api). Recorded from run 37241939097: model_parity rtlsim 32 (180 min), perf_gate rtlsim 32 (199), tensor_wg rtlsim 32 (185) and 64 (195) each split into two 90-100 min shards. Host cells. The planner now owns each cell's -m expression (`testcase.py expr`). A host cell selected ` and not ` with no driver term, so the regression host cell re-ran all 19 of the category's simx/rtlsim/xrtsim cases that their driver cells already run (~100 min per xlen). It now holds only driverless cases. Build reuse. Cases are ordered so those with one sim build signature run back to back: a blackbox case rebuilds the sim whenever its flags differ from the previous case's, so A, B, A compiled A twice (perf_gate rtlsim: 49 builds -> 32). ccache is re-enabled: it was disabled against the fmt::v8 stale-object class, which needs a cache populated against other headers -- impossible on a hosted runner, whose cache starts empty each job. Co-Authored-By: Claude Opus 5.5 --- .github/workflows/ci.yml | 36 +- ci/conftest.py | 52 ++ ci/record_runtimes.py | 102 ++++ ci/runtimes.json | 1216 ++++++++++++++++++++++++++++++++++++++ ci/testcase.py | 118 +++- 5 files changed, 1503 insertions(+), 21 deletions(-) create mode 100755 ci/record_runtimes.py create mode 100644 ci/runtimes.json diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 5942ad447a..a16733cf0e 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -23,8 +23,11 @@ on: description: "comma list: smoke,full,nightly (blank = all)" default: "" -env: - CCACHE_DISABLE: 1 +# ccache stays on: a hosted runner starts every job with an empty cache, so it +# cannot serve an object built against other headers (the stale-object class +# seen on persistent machines), and within a job it turns a cell's repeated sim +# builds and its config-invariant objects (Verilator runtime, host driver) into +# cache hits. jobs: # --------------------------------------------------------------------------- @@ -196,16 +199,22 @@ jobs: # --------------------------------------------------------------------------- # tests — one cell per (category, driver, xlen); pytest selects via markers. + # A cell whose recorded runtime exceeds ci/testcase.py's SHARD_BUDGET_S is + # planned as several shards (`shard` of `shards`), balanced on the per-case + # times in ci/runtimes.json, so every cell fits the 3 h timeout. # --------------------------------------------------------------------------- tests: needs: [plan, build] if: needs.plan.outputs.run == 'true' + name: tests (${{ matrix.category }}, ${{ matrix.driver }}, ${{ matrix.xlen }}${{ matrix.shards && format(', {0}/{1}', matrix.shard, matrix.shards) || '' }}) strategy: fail-fast: false matrix: include: ${{ fromJSON(needs.plan.outputs.cells) }} runs-on: ubuntu-22.04 - timeout-minutes: 300 + timeout-minutes: 180 + env: + CELL: ${{ matrix.category }}-${{ matrix.driver }}-${{ matrix.xlen }}${{ matrix.shards && format('-{0}of{1}', matrix.shard, matrix.shards) || '' }} steps: - uses: actions/checkout@v4 - name: Setup Vortex @@ -227,23 +236,18 @@ jobs: run: | cd build${{ matrix.xlen }} chmod -R +x . - EXPR="${{ matrix.category }}" - [ "${{ matrix.driver }}" != "host" ] && EXPR="$EXPR and ${{ matrix.driver }}" - # A cross-cutting check (model_parity, perf_gate) has its own dedicated - # cell whose `-m ` sweeps every such case catalog-wide; exclude - # those cases from every OTHER category's cell so each runs once, not - # twice. The check list comes from the planner, never a copy of it. - for c in $(python3 ci/testcase.py checks); do - [ "${{ matrix.category }}" = "$c" ] && continue - EXPR="$EXPR and not $c" - done + # The planner owns what a cell selects, so the cell runs exactly the + # cases it was planned (and, if sharded, split) from. + EXPR=$(python3 ci/testcase.py expr --cell="${{ matrix.category }}" --driver="${{ matrix.driver }}") + SHARD=() + [ -n "${{ matrix.shards }}" ] && SHARD=(--shard="${{ matrix.shard }}/${{ matrix.shards }}") VX_XLEN=${{ matrix.xlen }} python3 -m pytest ci -m "$EXPR" --strict-markers \ - -v \ - --junitxml="$GITHUB_WORKSPACE/junit-${{ matrix.category }}-${{ matrix.driver }}-${{ matrix.xlen }}.xml" + "${SHARD[@]}" -v \ + --junitxml="$GITHUB_WORKSPACE/junit-$CELL.xml" - uses: actions/upload-artifact@v4 if: always() with: - name: junit-${{ matrix.category }}-${{ matrix.driver }}-${{ matrix.xlen }} + name: junit-${{ env.CELL }} path: junit-*.xml complete: diff --git a/ci/conftest.py b/ci/conftest.py index 237794c84a..5c61785f06 100644 --- a/ci/conftest.py +++ b/ci/conftest.py @@ -37,6 +37,58 @@ def pytest_addoption(parser): parser.addoption("--update-baselines", action="store_true", default=False, help="perf_gate: record measured cycles into " "ci/baselines/perf/ instead of comparing") + parser.addoption("--shard", default=None, metavar="K/N", + help="run only shard K of N of the selected cases, split " + "on their recorded runtimes (ci/runtimes.json)") + + +@pytest.hookimpl(trylast=True) +def pytest_collection_modifyitems(config, items): + """Keep this --shard's slice of the selection, then run cases that compile + the same sim back to back. trylast: both act on what -m selected, which is + the set the planner sized the cell from.""" + _select_shard(config, items) + _group_by_sim(items) + + +def _group_by_sim(items): + """Order cases so those with one sim build signature run consecutively. + + A sim is rebuilt whenever a case's flags differ from the previous case's, so + a selection that alternates configs A, B, A compiles A twice. Each group + keeps the position of its first case, and cases keep their order within it. + """ + xlen = ambient_xlen() + first, keys = {}, [] + for pos, item in enumerate(items): + case = item.callspec.params.get("case") if hasattr(item, "callspec") else None + sig = case.sim_signature(xlen) if case is not None else None + key = sig if sig is not None else ("#", pos) # no sim: stays in place + first.setdefault(key, pos) + keys.append((first[key], pos)) + order = sorted(range(len(items)), key=lambda i: keys[i]) + items[:] = [items[i] for i in order] + + +def _select_shard(config, items): + spec = config.getoption("--shard") + if not spec: + return + try: + k, n = (int(x) for x in spec.split("/")) + except ValueError: + raise pytest.UsageError("--shard expects K/N, got {!r}".format(spec)) + if not 1 <= k <= n: + raise pytest.UsageError("--shard {}: K must be in 1..N".format(spec)) + ids = [item.callspec.id for item in items if hasattr(item, "callspec")] + _, plan = tc.shard_plan(ids, ambient_xlen(), n) + keep, drop = [], [] + for item in items: + cid = item.callspec.id if hasattr(item, "callspec") else None + (keep if plan.get(cid, 0) == k - 1 else drop).append(item) + if drop: + config.hook.pytest_deselected(items=drop) + items[:] = keep def pytest_sessionfinish(session): diff --git a/ci/record_runtimes.py b/ci/record_runtimes.py new file mode 100755 index 0000000000..dd2df3240c --- /dev/null +++ b/ci/record_runtimes.py @@ -0,0 +1,102 @@ +#!/usr/bin/env python3 + +# Copyright © 2019-2026 +# +# Licensed under the Apache License, Version 2.0 (the "License"). + +"""Record per-case CI runtimes into ci/runtimes.json from a GitHub Actions run. + +Reads the junit report every test cell of the run uploaded -- the wall time the +CI runners measured, never a local estimate -- and updates the per-xlen table +the planner splits long cells with (testcase.shard_plan). Cases the run did not +report keep their recorded time, so a run with a cancelled cell does not erase +that cell's history. + + ci/record_runtimes.py --run 37241939097 [--repo vortexgpgpu/vortex] + +Needs an authenticated `gh` CLI. +""" + +import argparse +import io +import json +import os +import re +import subprocess +import sys +import xml.etree.ElementTree as ET +import zipfile + +sys.path.insert(0, os.path.dirname(os.path.abspath(__file__))) +import testcase as tc # noqa: E402 + +# junit---[-of], as ci.yml's tests job names them. +ARTIFACT_RE = re.compile(r"^junit-.+-(32|64)(?:-\d+of\d+)?$") +CASE_RE = re.compile(r"^test_case\[(.+)\]$") + + +def gh_api(path, raw=False): + out = subprocess.run(["gh", "api", path], capture_output=True, check=True) + return out.stdout if raw else json.loads(out.stdout) + + +def artifacts(repo, run): + page, found = 1, [] + while True: + batch = gh_api("repos/{}/actions/runs/{}/artifacts?per_page=100&page={}" + .format(repo, run, page))["artifacts"] + found += batch + if len(batch) < 100: + return found + page += 1 + + +def case_times(blob): + """{case id: seconds} from one junit artifact zip.""" + times = {} + with zipfile.ZipFile(io.BytesIO(blob)) as zf: + for name in zf.namelist(): + if not name.endswith(".xml"): + continue + for case in ET.fromstring(zf.read(name)).iter("testcase"): + m = CASE_RE.match(case.get("name", "")) + if m and case.find("skipped") is None: + times[m.group(1)] = round(float(case.get("time", 0))) + return times + + +def main(): + p = argparse.ArgumentParser(description=__doc__.splitlines()[0]) + p.add_argument("--run", required=True, help="GitHub Actions run id of a CI run") + p.add_argument("--repo", default="vortexgpgpu/vortex") + args = p.parse_args() + + run = gh_api("repos/{}/actions/runs/{}".format(args.repo, args.run)) + try: + with open(tc.RUNTIMES_FILE) as fh: + table = json.load(fh) + except OSError: + table = {} + + updated = 0 + for art in artifacts(args.repo, args.run): + m = ARTIFACT_RE.match(art["name"]) + if not m or art.get("expired"): + continue + blob = gh_api("repos/{}/actions/artifacts/{}/zip".format(args.repo, art["id"]), + raw=True) + times = case_times(blob) + table.setdefault(m.group(1), {}).update(times) + updated += len(times) + + table["source"] = {"run": int(args.run), "sha": run["head_sha"]} + with open(tc.RUNTIMES_FILE, "w") as fh: + json.dump(table, fh, indent=1, sort_keys=True) + fh.write("\n") + print("recorded {} case times from run {} ({})".format( + updated, args.run, run["head_sha"][:9])) + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/ci/runtimes.json b/ci/runtimes.json new file mode 100644 index 0000000000..e886cb8119 --- /dev/null +++ b/ci/runtimes.json @@ -0,0 +1,1216 @@ +{ + "32": { + "amo:base:rtlsim": 36, + "amo:base:simx": 84, + "amo:isa-conformance:rtlsim": 55, + "amo:isa-conformance:simx": 100, + "amo:mc-l2:simx": 84, + "amo:mc-l3:rtlsim": 398, + "amo:mc-l3:simx": 84, + "amo:threads8:rtlsim": 36, + "amo:threads8:simx": 83, + "amo:wb-dirtybytes0:rtlsim": 35, + "amo:wb-dirtybytes1:rtlsim": 37, + "amo:wb:simx": 83, + "amo:zacas:rtlsim": 36, + "amo:zacas:simx": 83, + "cache:demo-1:rtlsim": 32, + "cache:demo-2:simx": 84, + "cache:demo-shared-icache-odd:simx": 82, + "cache:diverge-1:rtlsim": 191, + "cache:diverge-2:simx": 82, + "cache:diverge-3:rtlsim": 213, + "cache:diverge-4:simx": 82, + "cache:io_addr-1:rtlsim": 31, + "cache:io_addr-2:rtlsim": 32, + "cache:io_addr-3:simx": 81, + "cache:mstress-1:rtlsim": 29, + "cache:mstress-2:rtlsim": 29, + "cache:mstress-3:simx": 81, + "cache:mstress-4:rtlsim": 206, + "cache:mstress-5:simx": 82, + "cache:sgemm-10:rtlsim": 30, + "cache:sgemm-11:simx": 84, + "cache:sgemm-12:rtlsim": 37, + "cache:sgemm-13:simx": 83, + "cache:sgemm-14:rtlsim": 33, + "cache:sgemm-15:rtlsim": 33, + "cache:sgemm-16:rtlsim": 33, + "cache:sgemm-17:simx": 83, + "cache:sgemm-18:simx": 84, + "cache:sgemm-19:simx": 83, + "cache:sgemm-1:rtlsim": 34, + "cache:sgemm-20:rtlsim": 84, + "cache:sgemm-2:rtlsim": 34, + "cache:sgemm-3:rtlsim": 33, + "cache:sgemm-4:rtlsim": 33, + "cache:sgemm-5:rtlsim": 39, + "cache:sgemm-6:simx": 86, + "cache:sgemm-7:rtlsim": 33, + "cache:sgemm-8:rtlsim": 33, + "cache:sgemm-9:simx": 84, + "config1:demo-1:simx": 83, + "config1:demo-2:simx": 1, + "config1:demo-3:rtlsim": 31, + "config1:demo-4:rtlsim": 1, + "config1:diverge-10:simx": 82, + "config1:diverge-11:rtlsim": 91, + "config1:diverge-12:simx": 80, + "config1:diverge-13:rtlsim": 94, + "config1:diverge-14:simx": 80, + "config1:diverge-15:rtlsim": 67, + "config1:diverge-16:simx": 81, + "config1:diverge-17:rtlsim": 48, + "config1:diverge-18:rtlsim": 85, + "config1:diverge-19:simx": 84, + "config1:diverge-1:rtlsim": 33, + "config1:diverge-20:simx": 84, + "config1:diverge-21:rtlsim": 34, + "config1:diverge-22:rtlsim": 118, + "config1:diverge-23:rtlsim": 84, + "config1:diverge-24:simx": 83, + "config1:diverge-25:simx": 85, + "config1:diverge-26:simx": 84, + "config1:diverge-2:rtlsim": 23, + "config1:diverge-3:rtlsim": 24, + "config1:diverge-4:rtlsim": 25, + "config1:diverge-5:rtlsim": 36, + "config1:diverge-6:rtlsim": 31, + "config1:diverge-7:simx": 84, + "config1:diverge-8:simx": 85, + "config1:diverge-9:rtlsim": 91, + "config1:dogfood-1:rtlsim": 34, + "config1:dogfood-2:rtlsim": 36, + "config1:dogfood-3:simx": 86, + "config1:dogfood-4:simx": 85, + "config1:dogfood-5:rtlsim": 30, + "config1:dogfood-6:rtlsim": 31, + "config1:dogfood-7:rtlsim": 31, + "config1:dogfood-8:rtlsim": 31, + "config1:dogfood-9:rtlsim": 31, + "config1:regression/vecadd-1:rtlsim": 37, + "config1:regression/vecadd-2:rtlsim": 64, + "config1:regression/vecadd-3:simx": 83, + "config1:regression/vecadd-4:simx": 83, + "config1:regression/vecadd-5:rtlsim": 29, + "config1:regression/vecadd-6:rtlsim": 66, + "config1:regression/vecadd-7:simx": 82, + "config1:regression/vecadd-8:simx": 82, + "config1:sgemm-1:rtlsim": 84, + "config1:sgemm-2:rtlsim": 85, + "config2:demo-1:rtlsim": 16, + "config2:diverge-2:xrtsim": 5, + "config2:dogfood-2:xrtsim": 55, + "config2:dogfood-3:simx": 98, + "config2:dogfood-4:rtlsim": 26, + "config2:isa-32i-1:rtlsim": 25, + "config2:mstress-10:xrtsim": 19, + "config2:mstress-11:simx": 83, + "config2:mstress-12:simx": 83, + "config2:mstress-13:rtlsim": 16, + "config2:mstress-14:rtlsim": 20, + "config2:mstress-15:xrtsim": 23, + "config2:mstress-1:xrtsim": 18, + "config2:mstress-2:xrtsim": 20, + "config2:mstress-3:simx": 87, + "config2:mstress-4:rtlsim": 17, + "config2:mstress-5:simx": 83, + "config2:mstress-6:xrtsim": 18, + "config2:mstress-7:xrtsim": 19, + "config2:mstress-8:xrtsim": 19, + "config2:mstress-9:xrtsim": 19, + "config2:printf-1:xrtsim": 22, + "core:parity-raycast-nt16:rtlsim": 154, + "core:parity-raycast:rtlsim": 19, + "core:parity-sgemm-mc:rtlsim": 217, + "core:parity-sgemm:rtlsim": 10, + "core:parity-sgemmx:rtlsim": 146, + "core:parity-sgemv:rtlsim": 14, + "core:parity-softmax:rtlsim": 41, + "core:parity-stencil3d:rtlsim": 42, + "core:parity-vecadd:rtlsim": 114, + "core:raycast-nt16:rtlsim": 51, + "core:raycast-nt4:rtlsim": 35, + "core:sgemm-mc-nt16:rtlsim": 173, + "core:sgemm-mc-nt4:rtlsim": 107, + "core:sgemm-nt16:rtlsim": 65, + "core:sgemm-nt4:rtlsim": 56, + "core:sgemmx-nt16:rtlsim": 57, + "core:sgemmx-nt4:rtlsim": 45, + "core:sgemv-nt16:rtlsim": 47, + "core:sgemv-nt4:rtlsim": 29, + "core:softmax-nt16:rtlsim": 55, + "core:softmax-nt4:rtlsim": 47, + "core:stencil3d-nt16:rtlsim": 72, + "core:stencil3d-nt4:rtlsim": 45, + "core:vecadd-nt16:rtlsim": 41, + "core:vecadd-nt4:rtlsim": 26, + "debug:demo-1:opaesim": 55, + "debug:demo-2:xrtsim": 45, + "debug:demo-3:opaesim": 298, + "debug:demo-4:xrtsim": 137, + "debug:demo-5:simx": 63, + "debug:sgemm_tcu-1:simx": 74, + "debug:sgemm_tcu-2:rtlsim": 58, + "debug:simx_functional-1:simx": 81, + "debug:simx_functional-2:simx": 72, + "debug:simx_mt-1:simx": 74, + "dtm:legacy:None": 86, + "dxa:cta_cluster-1:rtlsim": 31, + "dxa:cta_cluster-1:simx": 89, + "dxa:cta_cluster-2c2s:rtlsim": 203, + "dxa:cta_cluster-2c2s:simx": 89, + "dxa:cta_cluster-mc:rtlsim": 65, + "dxa:cta_cluster-mc:simx": 88, + "dxa:dxa_copy-10:rtlsim": 1, + "dxa:dxa_copy-1:simx": 81, + "dxa:dxa_copy-2:simx": 87, + "dxa:dxa_copy-2c2s:rtlsim": 204, + "dxa:dxa_copy-2c2s:simx": 88, + "dxa:dxa_copy-3:simx": 87, + "dxa:dxa_copy-4:simx": 87, + "dxa:dxa_copy-5:simx": 1, + "dxa:dxa_copy-6:rtlsim": 30, + "dxa:dxa_copy-7:rtlsim": 33, + "dxa:dxa_copy-8:rtlsim": 51, + "dxa:dxa_copy-9:rtlsim": 32, + "dxa:dxa_copy_mcast-1:simx": 88, + "dxa:dxa_copy_mcast-2:rtlsim": 32, + "dxa:dxa_copy_mcast-2c2s:rtlsim": 206, + "dxa:dxa_copy_mcast-2c2s:simx": 88, + "dxa:dxa_copy_mcast-3:simx": 87, + "dxa:dxa_copy_mcast-4:rtlsim": 37, + "dxa:dxa_copy_mcast-5:simx": 88, + "dxa:dxa_copy_mcast-mc:rtlsim": 66, + "dxa:dxa_copy_mcast-mc:simx": 87, + "dxa:dxa_kmajor_check-1:simx": 91, + "dxa:dxa_kmajor_check-2:rtlsim": 31, + "dxa:model_parity-copy-mcast:rtlsim": 2, + "dxa:model_parity-copy:rtlsim": 125, + "dxa:model_parity-sgemm2-dxa-mcast:rtlsim": 133, + "dxa:model_parity-sgemm2-dxa:rtlsim": 133, + "dxa:model_parity-wgmma-dxa-mcast:rtlsim": 205, + "dxa:model_parity-wgmma-dxa:rtlsim": 3288, + "dxa:model_parity-wgmma-sp-dxa:rtlsim": 3465, + "dxa:perf_gate-copy-mcast:rtlsim": 26, + "dxa:perf_gate-copy-nt16:rtlsim": 42, + "dxa:perf_gate-copy-nt4:rtlsim": 26, + "dxa:perf_gate-sgemm2-dxa-mcast-nt16:rtlsim": 229, + "dxa:perf_gate-sgemm2-dxa-mcast-nt4:rtlsim": 96, + "dxa:perf_gate-sgemm2-dxa-nt16:rtlsim": 100, + "dxa:perf_gate-sgemm2-dxa-nt4:rtlsim": 101, + "dxa:perf_gate-wgmma-dxa-mcast-nt16:rtlsim": 1224, + "dxa:perf_gate-wgmma-dxa-mcast-nt4:rtlsim": 89, + "dxa:perf_gate-wgmma-dxa:rtlsim": 2541, + "dxa:perf_gate-wgmma-sp-dxa:rtlsim": 2652, + "dxa:sgemm2_dxa-10:rtlsim": 100, + "dxa:sgemm2_dxa-1:simx": 88, + "dxa:sgemm2_dxa-2:rtlsim": 32, + "dxa:sgemm2_dxa-3:simx": 88, + "dxa:sgemm2_dxa-4:rtlsim": 50, + "dxa:sgemm2_dxa-5:simx": 88, + "dxa:sgemm2_dxa-6:rtlsim": 49, + "dxa:sgemm2_dxa-7:simx": 88, + "dxa:sgemm2_dxa-8:rtlsim": 86, + "dxa:sgemm2_dxa-9:simx": 89, + "dxa:sgemm2_dxa_mcast-1:simx": 2, + "dxa:sgemm2_dxa_mcast-2:rtlsim": 2, + "dxa:sgemm2_dxa_mcast-3:simx": 1, + "dxa:sgemm2_dxa_mcast-mc:rtlsim": 66, + "dxa:sgemm2_dxa_mcast-mc:simx": 87, + "dxa:sgemm_tcu-1:simx": 94, + "dxa:sgemm_tcu-2:rtlsim": 36, + "dxa:sgemm_tcu_wg_dxa-1:simx": 95, + "dxa:sgemm_tcu_wg_dxa-2:rtlsim": 178, + "dxa:sgemm_tcu_wg_dxa-fedp2k:rtlsim": 831, + "dxa:sgemm_tcu_wg_dxa-fedp2k:simx": 96, + "dxa:sgemm_tcu_wg_dxa_mcast-1:simx": 95, + "dxa:sgemm_tcu_wg_dxa_mcast-2:rtlsim": 115, + "dxa:sgemm_tcu_wg_dxa_mcast-3:simx": 1, + "dxa:sgemm_tcu_wg_dxa_mcast-4:rtlsim": 6, + "dxa:sgemm_tcu_wg_dxa_mcast-mc:rtlsim": 396, + "dxa:sgemm_tcu_wg_dxa_mcast-mc:simx": 95, + "dxa:sgemm_tcu_wg_sp_dxa-1:simx": 100, + "dxa:sgemm_tcu_wg_sp_dxa-2:rtlsim": 211, + "gem5:legacy:None": 97, + "graphics:gfx_draw3d-11:rtlsim": 58, + "graphics:gfx_draw3d-12:simx": 82, + "graphics:gfx_draw3d-13:rtlsim": 192, + "graphics:gfx_draw3d-14:simx": 72, + "graphics:gfx_draw3d-15:rtlsim": 24, + "graphics:gfx_draw3d-16:simx": 80, + "graphics:gfx_draw3d-17:rtlsim": 29, + "graphics:gfx_draw3d-18:simx": 80, + "graphics:gfx_draw3d-19:rtlsim": 31, + "graphics:gfx_draw3d-1:simx": 83, + "graphics:gfx_draw3d-2:rtlsim": 72, + "graphics:gfx_draw3d-3:rtlsim": 66, + "graphics:gfx_draw3d-4:rtlsim": 185, + "graphics:gfx_draw3d-5:rtlsim": 198, + "graphics:gfx_draw3d-6:rtlsim": 64, + "graphics:gfx_draw3d-7:simx": 80, + "graphics:gfx_draw3d-9:xrtsim": 53, + "graphics:gfx_draw3d-sw_all:simx": 72, + "graphics:gfx_draw3d-sw_om:simx": 77, + "graphics:gfx_draw3d-sw_raster:simx": 77, + "graphics:gfx_draw3d-sw_tex:simx": 78, + "graphics:gfx_earlyz-box-rtlsim:rtlsim": 31, + "graphics:gfx_earlyz-box-simx:simx": 80, + "graphics:gfx_earlyz-evilskull128-rtlsim:rtlsim": 65, + "graphics:gfx_earlyz-evilskull32-rtlsim:rtlsim": 31, + "graphics:gfx_earlyz-evilskull32-simx:simx": 79, + "graphics:gfx_om-10:simx": 71, + "graphics:gfx_om-11:rtlsim": 47, + "graphics:gfx_om-12:rtlsim": 44, + "graphics:gfx_om-13:rtlsim": 84, + "graphics:gfx_om-1:simx": 70, + "graphics:gfx_om-2:simx": 2, + "graphics:gfx_om-3:simx": 2, + "graphics:gfx_om-4:simx": 2, + "graphics:gfx_om-5:rtlsim": 27, + "graphics:gfx_om-6:rtlsim": 5, + "graphics:gfx_om-7:rtlsim": 5, + "graphics:gfx_om-8:rtlsim": 5, + "graphics:gfx_om-9:rtlsim": 42, + "graphics:gfx_om-mrt-blend:rtlsim": 5, + "graphics:gfx_om-mrt-blend:simx": 2, + "graphics:gfx_om-mrt:rtlsim": 26, + "graphics:gfx_om-mrt:simx": 71, + "graphics:gfx_om-sw:rtlsim": 11, + "graphics:gfx_om-sw:simx": 4, + "graphics:gfx_pipeline_om:simx": 2, + "graphics:gfx_pipeline_raster:simx": 71, + "graphics:gfx_pipeline_tex:simx": 154, + "graphics:gfx_raster-1:simx": 70, + "graphics:gfx_raster-2:rtlsim": 20, + "graphics:gfx_raster-3:rtlsim": 23, + "graphics:gfx_raster-4:simx": 70, + "graphics:gfx_raster-5:simx": 70, + "graphics:gfx_raster-6:rtlsim": 19, + "graphics:gfx_raster-7:rtlsim": 20, + "graphics:gfx_raster-sw:simx": 72, + "graphics:gfx_tex-10:simx": 68, + "graphics:gfx_tex-11:rtlsim": 22, + "graphics:gfx_tex-12:simx": 68, + "graphics:gfx_tex-13:rtlsim": 22, + "graphics:gfx_tex-14:simx": 68, + "graphics:gfx_tex-15:rtlsim": 73, + "graphics:gfx_tex-16:rtlsim": 55, + "graphics:gfx_tex-17:rtlsim": 302, + "graphics:gfx_tex-1:rtlsim": 23, + "graphics:gfx_tex-2:rtlsim": 1, + "graphics:gfx_tex-3:rtlsim": 1, + "graphics:gfx_tex-4:rtlsim": 1, + "graphics:gfx_tex-5:rtlsim": 1, + "graphics:gfx_tex-6:rtlsim": 1, + "graphics:gfx_tex-7:rtlsim": 1, + "graphics:gfx_tex-8:simx": 69, + "graphics:gfx_tex-9:rtlsim": 1, + "graphics:gfx_tex-border-bilinear:rtlsim": 1, + "graphics:gfx_tex-border-bilinear:simx": 1, + "graphics:gfx_tex-border-point:rtlsim": 21, + "graphics:gfx_tex-border-point:simx": 70, + "graphics:gfx_tex4-ff:simx": 70, + "graphics:gfx_tex4-sw:simx": 3, + "graphics:model_parity-draw3d:rtlsim": 155, + "graphics:model_parity-om:rtlsim": 133, + "graphics:model_parity-raster:rtlsim": 124, + "graphics:model_parity-tex:rtlsim": 131, + "graphics:perf_gate-draw3d-mc:rtlsim": 129, + "graphics:perf_gate-draw3d-nt16:rtlsim": 67, + "graphics:perf_gate-draw3d-nt4:rtlsim": 41, + "graphics:perf_gate-om-nt16:rtlsim": 44, + "graphics:perf_gate-om-nt4:rtlsim": 32, + "graphics:perf_gate-raster-nt16:rtlsim": 40, + "graphics:perf_gate-raster-nt4:rtlsim": 27, + "graphics:perf_gate-tex-nt16:rtlsim": 51, + "graphics:perf_gate-tex-nt4:rtlsim": 28, + "graphics_parity:box-1core:rtlsim": 72, + "graphics_parity:box-1core:simx": 83, + "graphics_parity:box-2clu-2core:rtlsim": 359, + "graphics_parity:box-2clu-2core:simx": 84, + "graphics_parity:box-2core:rtlsim": 125, + "graphics_parity:box-2core:simx": 82, + "graphics_parity:box-4core-2raster:rtlsim": 334, + "graphics_parity:box-4core-2raster:simx": 83, + "graphics_parity:box-4core:rtlsim": 294, + "graphics_parity:box-4core:simx": 84, + "graphics_parity:triangle-2core:rtlsim": 83, + "graphics_parity:triangle-2core:simx": 79, + "hip:isa-1:simx": 132, + "hip:isa-4:xrtsim": 89, + "kernel:isa-1:simx": 54, + "kernel:isa-2:rtlsim": 22, + "mpi:isa-1:simx": 127, + "mpi:isa-2:simx": 1, + "mpi:isa-3:simx": 17, + "mpi:isa-4:simx": 1, + "mpi:isa-5:simx": 3, + "mpi:isa-6:simx": 3, + "mpi:isa-7:simx": 2, + "mpi:isa-8:simx": 2, + "opencl:hybridsort:rtlsim": 33, + "opencl:hybridsort:simx": 87, + "opencl:image-copy:rtlsim": 27, + "opencl:image-copy:simx": 84, + "opencl:image-ff-bgra:rtlsim": 6, + "opencl:image-ff-bgra:simx": 2, + "opencl:image-ff-linear:rtlsim": 6, + "opencl:image-ff-linear:simx": 2, + "opencl:image-ff:rtlsim": 333, + "opencl:image-ff:simx": 103, + "opencl:image-get:rtlsim": 2, + "opencl:image-get:simx": 2, + "opencl:image-hwsw:rtlsim": 412, + "opencl:image-hwsw:simx": 8, + "opencl:image-int:rtlsim": 2, + "opencl:image-int:simx": 2, + "opencl:image-linear:rtlsim": 3, + "opencl:image-linear:simx": 2, + "opencl:isa-1:simx": 308, + "opencl:isa-2:rtlsim": 544, + "opencl:lbm-1:simx": 91, + "opencl:lbm-2:rtlsim": 43, + "perf:sgemm-iw2:rtlsim": 108, + "perf:sgemm-iw2:simx": 85, + "perf:sgemm-iw4:rtlsim": 196, + "perf:sgemm-iw4:simx": 80, + "raytracing:model_parity-rt_raycast:rtlsim": 527, + "raytracing:perf_gate-rt_raycast-mc:rtlsim": 615, + "raytracing:perf_gate-rt_raycast-nt16:rtlsim": 416, + "raytracing:perf_gate-rt_raycast-nt4:rtlsim": 364, + "raytracing:rt_smoke_mc_a:simx": 65, + "raytracing:rt_smoke_mc_b:simx": 65, + "raytracing:rt_suite:rtlsim": 6805, + "raytracing:rt_suite:simx": 1405, + "regression:bar-slot-phase-1:simx": 64, + "regression:bar-slot-phase-2:xrtsim": 34, + "regression:dogfood-1:simx": 65, + "regression:dogfood-2:xrtsim": 57, + "regression:dogfood-3:simx": 65, + "regression:dogfood-4:xrtsim": 35, + "regression:gbar-phase-1:simx": 64, + "regression:gbar-phase-2:xrtsim": 55, + "regression:isa-1:simx": 1389, + "regression:isa-2:rtlsim": 1716, + "regression:model_parity-async-gbarrier:rtlsim": 131, + "regression:module_reload-1:simx": 0, + "regression:occupancy-1:simx": 66, + "regression:occupancy-2:rtlsim": 233, + "regression:occupancy-3:simx": 73, + "regression:occupancy-4:rtlsim": 1054, + "regression:reopen-1:xrtsim": 34, + "regression:reopen-2:None": 1, + "regression:reopen-3:None": 33, + "regression:single_cta-1:simx": 64, + "regression:single_cta-2:rtlsim": 27, + "regression:vecadd-1:simx": 72, + "riscv:isa-1:simx": 173, + "riscv:isa-2:rtlsim": 171, + "riscv:isa-32f-1:rtlsim": 0, + "riscv:isa-32f-2:rtlsim": 0, + "riscv:isa-32f-3:rtlsim": 0, + "riscv:isa-32f-4:rtlsim": 0, + "riscv:isa-32f-dsp:rtlsim": 21, + "riscv:isa-64d-1:rtlsim": 0, + "riscv:isa-64d-2:rtlsim": 0, + "riscv:isa-64f-1:rtlsim": 0, + "riscv:isa-64f-2:rtlsim": 0, + "riscv:isa-64f-3:rtlsim": 0, + "runtime:isa-1:None": 66, + "rvc:isa-1:simx": 86, + "rvc:isa-2:rtlsim": 33, + "rvc:isa-32c-1:simx": 100, + "rvc:isa-32c-2:rtlsim": 44, + "rvc:isa-3:simx": 3, + "rvc:isa-4:rtlsim": 6, + "rvc:isa-5:simx": 6, + "rvc:isa-6:rtlsim": 16, + "rvc:isa-7:simx": 9, + "rvc:isa-8:rtlsim": 17, + "scope:demo-1:opaesim": 72, + "scope:demo-2:xrtsim": 37, + "sst:legacy:None": 80, + "stress:axi-store-order:xrtsim": 73, + "stress:dogfood-1:opaesim": 332, + "stress:sgemm-1:xrtsim": 111, + "synthesis:base:None": 792, + "tensor:model_parity-fp16-mc:rtlsim": 239, + "tensor:model_parity-fp16:rtlsim": 153, + "tensor:model_parity-sgemm2-fp16:rtlsim": 265, + "tensor:perf_gate-fp16-isw4:rtlsim": 112, + "tensor:perf_gate-fp16-mc-nt16:rtlsim": 682, + "tensor:perf_gate-fp16-mc-nt4:rtlsim": 90, + "tensor:perf_gate-fp16:rtlsim": 50, + "tensor:perf_gate-sgemm2-fp16:rtlsim": 131, + "tensor:sgemm2_tcu-1:simx": 90, + "tensor:sgemm2_tcu-2:rtlsim": 128, + "tensor:sgemm_tcu-10:rtlsim": 27, + "tensor:sgemm_tcu-11:rtlsim": 27, + "tensor:sgemm_tcu-12:rtlsim": 27, + "tensor:sgemm_tcu-13:rtlsim": 27, + "tensor:sgemm_tcu-14:rtlsim": 26, + "tensor:sgemm_tcu-15:rtlsim": 26, + "tensor:sgemm_tcu-16:rtlsim": 27, + "tensor:sgemm_tcu-17:rtlsim": 27, + "tensor:sgemm_tcu-18:rtlsim": 45, + "tensor:sgemm_tcu-19:rtlsim": 270, + "tensor:sgemm_tcu-1:simx": 100, + "tensor:sgemm_tcu-20:rtlsim": 56, + "tensor:sgemm_tcu-21:rtlsim": 203, + "tensor:sgemm_tcu-2:simx": 89, + "tensor:sgemm_tcu-3:simx": 88, + "tensor:sgemm_tcu-4:simx": 88, + "tensor:sgemm_tcu-5:simx": 88, + "tensor:sgemm_tcu-6:simx": 88, + "tensor:sgemm_tcu-7:simx": 89, + "tensor:sgemm_tcu-8:rtlsim": 42, + "tensor:sgemm_tcu-9:rtlsim": 27, + "tensor:sgemm_tcu-fedp2k:simx": 89, + "tensor_mx:model_parity-mxfp8:rtlsim": 129, + "tensor_mx:perf_gate-mxfp8:rtlsim": 29, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt16:rtlsim": 36, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt32:rtlsim": 51, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt4:rtlsim": 18, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt8:rtlsim": 21, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt16:rtlsim": 35, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt32:rtlsim": 54, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt4:rtlsim": 21, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt8:rtlsim": 22, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt16:rtlsim": 35, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt32:rtlsim": 51, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt4:rtlsim": 20, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt8:rtlsim": 21, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt16:rtlsim": 36, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt32:rtlsim": 55, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt4-tls:rtlsim": 20, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt4:rtlsim": 20, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt8:rtlsim": 21, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt16:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt32:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt4:simx": 49, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt8:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt16:simx": 51, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt32:simx": 47, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt4:simx": 49, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt8:simx": 49, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt16:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt32:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt4:simx": 49, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt8:simx": 50, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt16:simx": 47, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt32:simx": 49, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt4-tls:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt4:simx": 48, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt8:simx": 47, + "tensor_sp:model_parity-sparse-fp16:rtlsim": 152, + "tensor_sp:perf_gate-sparse-fp16:rtlsim": 50, + "tensor_sp:sgemm_tcu_sp-10:simx": 71, + "tensor_sp:sgemm_tcu_sp-11:simx": 72, + "tensor_sp:sgemm_tcu_sp-12:simx": 71, + "tensor_sp:sgemm_tcu_sp-13:rtlsim": 32, + "tensor_sp:sgemm_tcu_sp-14:rtlsim": 31, + "tensor_sp:sgemm_tcu_sp-15:rtlsim": 32, + "tensor_sp:sgemm_tcu_sp-16:rtlsim": 33, + "tensor_sp:sgemm_tcu_sp-17:simx": 69, + "tensor_sp:sgemm_tcu_sp-18:simx": 69, + "tensor_sp:sgemm_tcu_sp-19:simx": 71, + "tensor_sp:sgemm_tcu_sp-1:simx": 73, + "tensor_sp:sgemm_tcu_sp-20:simx": 71, + "tensor_sp:sgemm_tcu_sp-21:rtlsim": 37, + "tensor_sp:sgemm_tcu_sp-22:rtlsim": 38, + "tensor_sp:sgemm_tcu_sp-23:rtlsim": 37, + "tensor_sp:sgemm_tcu_sp-24:rtlsim": 37, + "tensor_sp:sgemm_tcu_sp-25:simx": 71, + "tensor_sp:sgemm_tcu_sp-26:simx": 71, + "tensor_sp:sgemm_tcu_sp-27:simx": 71, + "tensor_sp:sgemm_tcu_sp-28:simx": 71, + "tensor_sp:sgemm_tcu_sp-29:rtlsim": 56, + "tensor_sp:sgemm_tcu_sp-2:simx": 69, + "tensor_sp:sgemm_tcu_sp-30:rtlsim": 55, + "tensor_sp:sgemm_tcu_sp-31:rtlsim": 56, + "tensor_sp:sgemm_tcu_sp-32:rtlsim": 58, + "tensor_sp:sgemm_tcu_sp-33:simx": 71, + "tensor_sp:sgemm_tcu_sp-34:simx": 71, + "tensor_sp:sgemm_tcu_sp-35:simx": 70, + "tensor_sp:sgemm_tcu_sp-36:simx": 72, + "tensor_sp:sgemm_tcu_sp-37:rtlsim": 83, + "tensor_sp:sgemm_tcu_sp-38:rtlsim": 80, + "tensor_sp:sgemm_tcu_sp-39:rtlsim": 78, + "tensor_sp:sgemm_tcu_sp-3:simx": 71, + "tensor_sp:sgemm_tcu_sp-40:rtlsim": 81, + "tensor_sp:sgemm_tcu_sp-4:simx": 71, + "tensor_sp:sgemm_tcu_sp-5:rtlsim": 59, + "tensor_sp:sgemm_tcu_sp-6:rtlsim": 29, + "tensor_sp:sgemm_tcu_sp-7:rtlsim": 29, + "tensor_sp:sgemm_tcu_sp-8:rtlsim": 29, + "tensor_sp:sgemm_tcu_sp-9:simx": 70, + "tensor_sp_mx:sgemm_tcu_sp_mx-rtlsim-mxfp8-nt4:rtlsim": 40, + "tensor_sp_mx:sgemm_tcu_sp_mx-rtlsim-mxfp8-nt8:rtlsim": 37, + "tensor_sp_mx:sgemm_tcu_sp_mx-simx-mxfp8-nt4:simx": 70, + "tensor_sp_mx:sgemm_tcu_sp_mx-simx-mxfp8-nt8:simx": 70, + "tensor_wg:model_parity-wgmma-fedp2k-rs:rtlsim": 160, + "tensor_wg:model_parity-wgmma-int8:rtlsim": 235, + "tensor_wg:perf_gate-wgmma-fedp2k-rs:rtlsim": 608, + "tensor_wg:perf_gate-wgmma-fp16-ss:rtlsim": 144, + "tensor_wg:perf_gate-wgmma-sparse:rtlsim": 130, + "tensor_wg:sgemm_tcu_wg-10:simx": 85, + "tensor_wg:sgemm_tcu_wg-1:simx": 87, + "tensor_wg:sgemm_tcu_wg-2:simx": 85, + "tensor_wg:sgemm_tcu_wg-3:simx": 85, + "tensor_wg:sgemm_tcu_wg-4:simx": 85, + "tensor_wg:sgemm_tcu_wg-5:rtlsim": 145, + "tensor_wg:sgemm_tcu_wg-6:rtlsim": 1606, + "tensor_wg:sgemm_tcu_wg-7:rtlsim": 3993, + "tensor_wg:sgemm_tcu_wg-8:simx": 84, + "tensor_wg:sgemm_tcu_wg-9:rtlsim": 142, + "tensor_wg:sgemm_tcu_wg-fedp2k:simx": 85, + "tensor_wg:sgemm_tcu_wg_mx-baseline:simx": 85, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-nvfp4:simx": 86, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-rs:simx": 85, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-ss:simx": 86, + "tensor_wg:sgemm_tcu_wg_sp-1:simx": 86, + "tensor_wg:sgemm_tcu_wg_sp-2:simx": 86, + "tensor_wg:sgemm_tcu_wg_sp-3:simx": 85, + "tensor_wg:sgemm_tcu_wg_sp-4:rtlsim": 151, + "tensor_wg:sgemm_tcu_wg_sp-5:rtlsim": 1628, + "tensor_wg:sgemm_tcu_wg_sp-6:rtlsim": 3305, + "tensor_wg:sgemm_tcu_wg_sp-7:simx": 86, + "tensor_wg:sgemm_tcu_wg_sp-8:rtlsim": 149, + "tensor_wg:sgemm_tcu_wg_sp-fedp2k-rs:simx": 85, + "tensor_wg:sgemm_tcu_wg_sp-fedp2k-ss:simx": 86, + "unittest:hw-fpu:None": 2, + "unittest:hw-tcu-dsp:None": 71, + "unittest:hw-tcu:None": 222, + "unittest:hw-tex-lerp:None": 1, + "unittest:hw:None": 254, + "unittest:sw:None": 183, + "vm:isa-10:rtlsim": 48, + "vm:isa-10:simx": 72, + "vm:isa-1:rtlsim": 63, + "vm:isa-1:simx": 113, + "vm:isa-2:rtlsim": 7, + "vm:isa-2:simx": 2, + "vm:isa-3:rtlsim": 19, + "vm:isa-3:simx": 4, + "vm:isa-4:rtlsim": 16, + "vm:isa-4:simx": 5, + "vm:isa-5:rtlsim": 48, + "vm:isa-5:simx": 70, + "vm:isa-6:rtlsim": 61, + "vm:isa-6:simx": 118, + "vm:isa-7:rtlsim": 6, + "vm:isa-7:simx": 2, + "vm:isa-8:rtlsim": 19, + "vm:isa-8:simx": 4, + "vm:isa-9:rtlsim": 16, + "vm:isa-9:simx": 5, + "vm:mc-l2-amo:rtlsim": 1369, + "vm:mc-l2-amo:simx": 121, + "vm:mc-l3:rtlsim": 503, + "vm:mc-l3:simx": 122, + "vm:parity-vm-sgemm:rtlsim": 137, + "vm:parity-vm-vecadd:rtlsim": 5, + "vm:tlb-concurrent:rtlsim": 617, + "vm:tlb-concurrent:simx": 112, + "vm:tlb-stress-l1:rtlsim": 29, + "vm:tlb-stress-l1:simx": 59, + "vm:tlb-stress-l2:rtlsim": 1, + "vm:tlb-stress-l2:simx": 0, + "vm:tlb-stress-walk:rtlsim": 1, + "vm:tlb-stress-walk:simx": 0, + "vm:vm-amo:rtlsim": 60, + "vm:vm-amo:simx": 115, + "vm:vm-drain:rtlsim": 0, + "vm:vm-drain:simx": 0, + "vm:vm-fault:rtlsim": 28, + "vm:vm-fault:simx": 58, + "vm:vm-fence:rtlsim": 0, + "vm:vm-fence:simx": 0, + "vm:vm-stress:rtlsim": 1, + "vm:vm-stress:simx": 1, + "vm:vm-superpage:rtlsim": 29, + "vm:vm-superpage:simx": 55, + "vulkan:ff-all-sw:simx": 88, + "vulkan:ff-om-only:simx": 93, + "vulkan:ff-raster-om:simx": 103, + "vulkan:ff-raster-only:simx": 91, + "vulkan:ff-raster-tex-om:simx": 101, + "vulkan:ff-raster-tex:simx": 97, + "vulkan:ff-tex-om:simx": 109, + "vulkan:ff-tex-only:simx": 94, + "vulkan:gfx-multidraw:rtlsim": 52, + "vulkan:gfx-multidraw:simx": 102, + "vulkan:gfx-rt-rtquery-id:simx": 115, + "vulkan:isa-1:simx": 1749, + "vulkan:isa-4:xrtsim": 6496, + "vulkan:rt-raytrace:rtlsim": 726, + "vulkan:rt-raytrace:simx": 114, + "vulkan:rt-rtquery:simx": 114 + }, + "64": { + "amo:base:rtlsim": 38, + "amo:base:simx": 66, + "amo:isa-conformance:rtlsim": 70, + "amo:isa-conformance:simx": 95, + "amo:mc-l2:simx": 67, + "amo:mc-l3:rtlsim": 577, + "amo:mc-l3:simx": 68, + "amo:threads8:rtlsim": 40, + "amo:threads8:simx": 66, + "amo:wb-dirtybytes0:rtlsim": 38, + "amo:wb-dirtybytes1:rtlsim": 37, + "amo:wb:simx": 66, + "amo:zacas:rtlsim": 39, + "amo:zacas:simx": 66, + "cache:demo-1:rtlsim": 26, + "cache:demo-2:simx": 66, + "cache:demo-shared-icache-odd:simx": 64, + "cache:diverge-1:rtlsim": 181, + "cache:diverge-2:simx": 67, + "cache:diverge-3:rtlsim": 176, + "cache:diverge-4:simx": 67, + "cache:io_addr-1:rtlsim": 28, + "cache:io_addr-2:rtlsim": 28, + "cache:io_addr-3:simx": 63, + "cache:mstress-1:rtlsim": 24, + "cache:mstress-2:rtlsim": 25, + "cache:mstress-3:simx": 66, + "cache:mstress-4:rtlsim": 178, + "cache:mstress-5:simx": 66, + "cache:sgemm-10:rtlsim": 30, + "cache:sgemm-11:simx": 70, + "cache:sgemm-12:rtlsim": 38, + "cache:sgemm-13:simx": 69, + "cache:sgemm-14:rtlsim": 34, + "cache:sgemm-15:rtlsim": 33, + "cache:sgemm-16:rtlsim": 34, + "cache:sgemm-17:simx": 70, + "cache:sgemm-18:simx": 67, + "cache:sgemm-19:simx": 67, + "cache:sgemm-1:rtlsim": 34, + "cache:sgemm-20:rtlsim": 111, + "cache:sgemm-2:rtlsim": 34, + "cache:sgemm-3:rtlsim": 34, + "cache:sgemm-4:rtlsim": 34, + "cache:sgemm-5:rtlsim": 38, + "cache:sgemm-6:simx": 68, + "cache:sgemm-7:rtlsim": 33, + "cache:sgemm-8:rtlsim": 34, + "cache:sgemm-9:simx": 67, + "config1:demo-1:simx": 65, + "config1:demo-2:simx": 0, + "config1:demo-3:rtlsim": 36, + "config1:demo-4:rtlsim": 1, + "config1:diverge-10:simx": 65, + "config1:diverge-11:rtlsim": 128, + "config1:diverge-12:simx": 65, + "config1:diverge-13:rtlsim": 131, + "config1:diverge-14:simx": 65, + "config1:diverge-15:rtlsim": 78, + "config1:diverge-16:simx": 64, + "config1:diverge-17:rtlsim": 58, + "config1:diverge-18:rtlsim": 108, + "config1:diverge-19:simx": 65, + "config1:diverge-1:rtlsim": 26, + "config1:diverge-20:simx": 66, + "config1:diverge-21:rtlsim": 39, + "config1:diverge-22:rtlsim": 106, + "config1:diverge-23:rtlsim": 109, + "config1:diverge-24:simx": 65, + "config1:diverge-25:simx": 67, + "config1:diverge-26:simx": 65, + "config1:diverge-2:rtlsim": 25, + "config1:diverge-3:rtlsim": 28, + "config1:diverge-4:rtlsim": 28, + "config1:diverge-5:rtlsim": 44, + "config1:diverge-6:rtlsim": 34, + "config1:diverge-7:simx": 70, + "config1:diverge-8:simx": 66, + "config1:diverge-9:rtlsim": 119, + "config1:dogfood-1:rtlsim": 27, + "config1:dogfood-2:rtlsim": 29, + "config1:dogfood-3:simx": 65, + "config1:dogfood-4:simx": 65, + "config1:dogfood-5:rtlsim": 34, + "config1:dogfood-6:rtlsim": 34, + "config1:dogfood-7:rtlsim": 34, + "config1:dogfood-8:rtlsim": 34, + "config1:dogfood-9:rtlsim": 34, + "config1:regression/vecadd-1:rtlsim": 45, + "config1:regression/vecadd-2:rtlsim": 80, + "config1:regression/vecadd-3:simx": 66, + "config1:regression/vecadd-4:simx": 65, + "config1:regression/vecadd-5:rtlsim": 36, + "config1:regression/vecadd-6:rtlsim": 80, + "config1:regression/vecadd-7:simx": 64, + "config1:regression/vecadd-8:simx": 64, + "config1:sgemm-1:rtlsim": 125, + "config1:sgemm-2:rtlsim": 125, + "config2:demo-1:rtlsim": 31, + "config2:diverge-2:xrtsim": 10, + "config2:dogfood-2:xrtsim": 61, + "config2:dogfood-3:simx": 52, + "config2:dogfood-4:rtlsim": 38, + "config2:isa-32i-1:rtlsim": 64, + "config2:mstress-10:xrtsim": 39, + "config2:mstress-11:simx": 58, + "config2:mstress-12:simx": 57, + "config2:mstress-13:rtlsim": 37, + "config2:mstress-14:rtlsim": 49, + "config2:mstress-15:xrtsim": 63, + "config2:mstress-1:xrtsim": 38, + "config2:mstress-2:xrtsim": 43, + "config2:mstress-3:simx": 54, + "config2:mstress-4:rtlsim": 41, + "config2:mstress-5:simx": 56, + "config2:mstress-6:xrtsim": 39, + "config2:mstress-7:xrtsim": 40, + "config2:mstress-8b:xrtsim": 1, + "config2:mstress-9:xrtsim": 40, + "config2:printf-1:xrtsim": 53, + "cupbop:legacy:None": 100, + "debug:demo-1:opaesim": 56, + "debug:demo-2:xrtsim": 37, + "debug:demo-3:opaesim": 351, + "debug:demo-4:xrtsim": 127, + "debug:demo-5:simx": 77, + "debug:sgemm_tcu-1:simx": 94, + "debug:sgemm_tcu-2:rtlsim": 70, + "debug:simx_functional-1:simx": 115, + "debug:simx_functional-2:simx": 103, + "debug:simx_mt-1:simx": 104, + "dtm:legacy:None": 87, + "dxa:cta_cluster-1:rtlsim": 25, + "dxa:cta_cluster-1:simx": 80, + "dxa:cta_cluster-2c2s:rtlsim": 124, + "dxa:cta_cluster-2c2s:simx": 79, + "dxa:cta_cluster-mc:rtlsim": 55, + "dxa:cta_cluster-mc:simx": 79, + "dxa:dxa_copy-10:rtlsim": 1, + "dxa:dxa_copy-1:simx": 74, + "dxa:dxa_copy-2:simx": 79, + "dxa:dxa_copy-2c2s:rtlsim": 125, + "dxa:dxa_copy-2c2s:simx": 80, + "dxa:dxa_copy-3:simx": 79, + "dxa:dxa_copy-4:simx": 79, + "dxa:dxa_copy-5:simx": 1, + "dxa:dxa_copy-6:rtlsim": 23, + "dxa:dxa_copy-7:rtlsim": 26, + "dxa:dxa_copy-8:rtlsim": 45, + "dxa:dxa_copy-9:rtlsim": 25, + "dxa:dxa_copy_mcast-1:simx": 80, + "dxa:dxa_copy_mcast-2:rtlsim": 25, + "dxa:dxa_copy_mcast-2c2s:rtlsim": 124, + "dxa:dxa_copy_mcast-2c2s:simx": 80, + "dxa:dxa_copy_mcast-3:simx": 80, + "dxa:dxa_copy_mcast-4:rtlsim": 29, + "dxa:dxa_copy_mcast-5:simx": 81, + "dxa:dxa_copy_mcast-mc:rtlsim": 55, + "dxa:dxa_copy_mcast-mc:simx": 80, + "dxa:dxa_kmajor_check-1:simx": 81, + "dxa:dxa_kmajor_check-2:rtlsim": 26, + "dxa:sgemm2_dxa-10:rtlsim": 99, + "dxa:sgemm2_dxa-1:simx": 79, + "dxa:sgemm2_dxa-2:rtlsim": 26, + "dxa:sgemm2_dxa-3:simx": 80, + "dxa:sgemm2_dxa-4:rtlsim": 46, + "dxa:sgemm2_dxa-5:simx": 80, + "dxa:sgemm2_dxa-6:rtlsim": 46, + "dxa:sgemm2_dxa-7:simx": 80, + "dxa:sgemm2_dxa-8:rtlsim": 72, + "dxa:sgemm2_dxa-9:simx": 80, + "dxa:sgemm2_dxa_mcast-1:simx": 1, + "dxa:sgemm2_dxa_mcast-2:rtlsim": 2, + "dxa:sgemm2_dxa_mcast-3:simx": 1, + "dxa:sgemm2_dxa_mcast-mc:rtlsim": 55, + "dxa:sgemm2_dxa_mcast-mc:simx": 79, + "dxa:sgemm_tcu-1:simx": 85, + "dxa:sgemm_tcu-2:rtlsim": 31, + "dxa:sgemm_tcu_wg_dxa-1:simx": 86, + "dxa:sgemm_tcu_wg_dxa-2:rtlsim": 207, + "dxa:sgemm_tcu_wg_dxa-fedp2k:rtlsim": 648, + "dxa:sgemm_tcu_wg_dxa-fedp2k:simx": 86, + "dxa:sgemm_tcu_wg_dxa_mcast-1:simx": 86, + "dxa:sgemm_tcu_wg_dxa_mcast-2:rtlsim": 125, + "dxa:sgemm_tcu_wg_dxa_mcast-3:simx": 1, + "dxa:sgemm_tcu_wg_dxa_mcast-4:rtlsim": 6, + "dxa:sgemm_tcu_wg_dxa_mcast-mc:rtlsim": 297, + "dxa:sgemm_tcu_wg_dxa_mcast-mc:simx": 86, + "dxa:sgemm_tcu_wg_sp_dxa-1:simx": 87, + "dxa:sgemm_tcu_wg_sp_dxa-2:rtlsim": 204, + "gem5:legacy:None": 74, + "graphics:gfx_draw3d-11:rtlsim": 87, + "graphics:gfx_draw3d-12:simx": 107, + "graphics:gfx_draw3d-13:rtlsim": 322, + "graphics:gfx_draw3d-14:simx": 94, + "graphics:gfx_draw3d-15:rtlsim": 32, + "graphics:gfx_draw3d-16:simx": 101, + "graphics:gfx_draw3d-17:rtlsim": 38, + "graphics:gfx_draw3d-18:simx": 101, + "graphics:gfx_draw3d-19:rtlsim": 42, + "graphics:gfx_draw3d-1:simx": 106, + "graphics:gfx_draw3d-2:rtlsim": 109, + "graphics:gfx_draw3d-3:rtlsim": 103, + "graphics:gfx_draw3d-4:rtlsim": 280, + "graphics:gfx_draw3d-5:rtlsim": 320, + "graphics:gfx_draw3d-6:rtlsim": 85, + "graphics:gfx_draw3d-7:simx": 103, + "graphics:gfx_draw3d-9:xrtsim": 44, + "graphics:gfx_draw3d-sw_all:simx": 93, + "graphics:gfx_draw3d-sw_om:simx": 99, + "graphics:gfx_draw3d-sw_raster:simx": 100, + "graphics:gfx_draw3d-sw_tex:simx": 100, + "graphics:gfx_earlyz-box-rtlsim:rtlsim": 43, + "graphics:gfx_earlyz-box-simx:simx": 102, + "graphics:gfx_earlyz-evilskull128-rtlsim:rtlsim": 82, + "graphics:gfx_earlyz-evilskull32-rtlsim:rtlsim": 42, + "graphics:gfx_earlyz-evilskull32-simx:simx": 102, + "graphics:gfx_om-10:simx": 92, + "graphics:gfx_om-11:rtlsim": 64, + "graphics:gfx_om-12:rtlsim": 60, + "graphics:gfx_om-13:rtlsim": 138, + "graphics:gfx_om-1:simx": 91, + "graphics:gfx_om-2:simx": 3, + "graphics:gfx_om-3:simx": 3, + "graphics:gfx_om-4:simx": 3, + "graphics:gfx_om-5:rtlsim": 33, + "graphics:gfx_om-6:rtlsim": 6, + "graphics:gfx_om-7:rtlsim": 6, + "graphics:gfx_om-8:rtlsim": 6, + "graphics:gfx_om-9:rtlsim": 55, + "graphics:gfx_om-mrt-blend:rtlsim": 7, + "graphics:gfx_om-mrt-blend:simx": 3, + "graphics:gfx_om-mrt:rtlsim": 34, + "graphics:gfx_om-mrt:simx": 92, + "graphics:gfx_om-sw:rtlsim": 15, + "graphics:gfx_om-sw:simx": 5, + "graphics:gfx_pipeline_om:simx": 3, + "graphics:gfx_pipeline_raster:simx": 90, + "graphics:gfx_pipeline_tex:simx": 200, + "graphics:gfx_raster-1:simx": 90, + "graphics:gfx_raster-2:rtlsim": 29, + "graphics:gfx_raster-3:rtlsim": 34, + "graphics:gfx_raster-4:simx": 89, + "graphics:gfx_raster-5:simx": 90, + "graphics:gfx_raster-6:rtlsim": 28, + "graphics:gfx_raster-7:rtlsim": 28, + "graphics:gfx_raster-sw:simx": 92, + "graphics:gfx_tex-10:simx": 89, + "graphics:gfx_tex-11:rtlsim": 33, + "graphics:gfx_tex-12:simx": 89, + "graphics:gfx_tex-13:rtlsim": 32, + "graphics:gfx_tex-14:simx": 89, + "graphics:gfx_tex-15:rtlsim": 107, + "graphics:gfx_tex-16:rtlsim": 71, + "graphics:gfx_tex-17:rtlsim": 483, + "graphics:gfx_tex-1:rtlsim": 33, + "graphics:gfx_tex-2:rtlsim": 2, + "graphics:gfx_tex-3:rtlsim": 2, + "graphics:gfx_tex-4:rtlsim": 2, + "graphics:gfx_tex-5:rtlsim": 2, + "graphics:gfx_tex-6:rtlsim": 2, + "graphics:gfx_tex-7:rtlsim": 2, + "graphics:gfx_tex-8:simx": 91, + "graphics:gfx_tex-9:rtlsim": 2, + "graphics:gfx_tex-border-bilinear:rtlsim": 1, + "graphics:gfx_tex-border-bilinear:simx": 1, + "graphics:gfx_tex-border-point:rtlsim": 31, + "graphics:gfx_tex-border-point:simx": 89, + "graphics:gfx_tex4-ff:simx": 89, + "graphics:gfx_tex4-sw:simx": 3, + "graphics_parity:box-1core:rtlsim": 84, + "graphics_parity:box-1core:simx": 107, + "graphics_parity:box-2clu-2core:rtlsim": 403, + "graphics_parity:box-2clu-2core:simx": 117, + "graphics_parity:box-2core:rtlsim": 155, + "graphics_parity:box-2core:simx": 112, + "graphics_parity:box-4core-2raster:rtlsim": 364, + "graphics_parity:box-4core-2raster:simx": 114, + "graphics_parity:box-4core:rtlsim": 348, + "graphics_parity:box-4core:simx": 113, + "graphics_parity:triangle-2core:rtlsim": 97, + "graphics_parity:triangle-2core:simx": 109, + "hip:isa-1:simx": 178, + "hip:isa-4:xrtsim": 101, + "kernel:isa-1:simx": 88, + "kernel:isa-2:rtlsim": 49, + "mpi:isa-1:simx": 130, + "mpi:isa-2:simx": 1, + "mpi:isa-3:simx": 17, + "mpi:isa-4:simx": 1, + "mpi:isa-5:simx": 5, + "mpi:isa-6:simx": 4, + "mpi:isa-7:simx": 2, + "mpi:isa-8:simx": 2, + "opencl:image-copy:rtlsim": 3, + "opencl:image-copy:simx": 2, + "opencl:image-ff-bgra:rtlsim": 3, + "opencl:image-ff-bgra:simx": 2, + "opencl:image-ff-linear:rtlsim": 3, + "opencl:image-ff-linear:simx": 2, + "opencl:image-ff:rtlsim": 67, + "opencl:image-ff:simx": 97, + "opencl:image-get:rtlsim": 2, + "opencl:image-get:simx": 2, + "opencl:image-hwsw:rtlsim": 108, + "opencl:image-hwsw:simx": 7, + "opencl:image-int:rtlsim": 2, + "opencl:image-int:simx": 2, + "opencl:image-linear:rtlsim": 3, + "opencl:image-linear:simx": 2, + "opencl:isa-1:simx": 251, + "opencl:isa-2:rtlsim": 367, + "opencl:lbm-1:simx": 84, + "opencl:lbm-2:rtlsim": 57, + "regression:bar-slot-phase-1:simx": 47, + "regression:bar-slot-phase-2:xrtsim": 22, + "regression:dogfood-1:simx": 47, + "regression:dogfood-2:xrtsim": 38, + "regression:dogfood-3:simx": 46, + "regression:dogfood-4:xrtsim": 22, + "regression:gbar-phase-1:simx": 45, + "regression:gbar-phase-2:xrtsim": 37, + "regression:isa-1:simx": 951, + "regression:isa-2:rtlsim": 1980, + "regression:module_reload-1:simx": 0, + "regression:occupancy-1:simx": 48, + "regression:occupancy-2:rtlsim": 317, + "regression:occupancy-3:simx": 53, + "regression:occupancy-4:rtlsim": 1473, + "regression:reopen-1:xrtsim": 21, + "regression:reopen-2:None": 1, + "regression:reopen-3:None": 37, + "regression:single_cta-1:simx": 46, + "regression:single_cta-2:rtlsim": 31, + "regression:vecadd-1:simx": 52, + "riscv:isa-1:simx": 152, + "riscv:isa-2:rtlsim": 158, + "riscv:isa-32f-1:rtlsim": 0, + "riscv:isa-32f-2:rtlsim": 0, + "riscv:isa-32f-3:rtlsim": 0, + "riscv:isa-32f-4:rtlsim": 0, + "riscv:isa-64d-1:rtlsim": 0, + "riscv:isa-64d-2:rtlsim": 0, + "riscv:isa-64f-1:rtlsim": 0, + "riscv:isa-64f-2:rtlsim": 0, + "riscv:isa-64f-3:rtlsim": 0, + "runtime:isa-1:None": 61, + "rvc:isa-1:simx": 86, + "rvc:isa-2:rtlsim": 47, + "rvc:isa-32c-1:simx": 116, + "rvc:isa-32c-2:rtlsim": 76, + "rvc:isa-3:simx": 3, + "rvc:isa-4:rtlsim": 7, + "rvc:isa-5:simx": 3, + "rvc:isa-64c-1:simx": 0, + "rvc:isa-64c-2:rtlsim": 0, + "rvc:isa-6:rtlsim": 5, + "rvc:isa-7:simx": 9, + "rvc:isa-8:rtlsim": 19, + "scope:demo-1:opaesim": 43, + "scope:demo-2:xrtsim": 54, + "sst:legacy:None": 80, + "stress:axi-store-order:xrtsim": 74, + "stress:dogfood-1:opaesim": 658, + "stress:sgemm-1:xrtsim": 178, + "synthesis:base:None": 2100, + "tensor:sgemm2_tcu-1:simx": 90, + "tensor:sgemm2_tcu-2:rtlsim": 104, + "tensor:sgemm_tcu-10:rtlsim": 27, + "tensor:sgemm_tcu-11:rtlsim": 27, + "tensor:sgemm_tcu-12:rtlsim": 27, + "tensor:sgemm_tcu-13:rtlsim": 28, + "tensor:sgemm_tcu-14:rtlsim": 28, + "tensor:sgemm_tcu-15:rtlsim": 28, + "tensor:sgemm_tcu-16:rtlsim": 29, + "tensor:sgemm_tcu-17:rtlsim": 29, + "tensor:sgemm_tcu-18:rtlsim": 48, + "tensor:sgemm_tcu-19:rtlsim": 256, + "tensor:sgemm_tcu-1:simx": 100, + "tensor:sgemm_tcu-20:rtlsim": 55, + "tensor:sgemm_tcu-21:rtlsim": 186, + "tensor:sgemm_tcu-2:simx": 89, + "tensor:sgemm_tcu-3:simx": 90, + "tensor:sgemm_tcu-4:simx": 91, + "tensor:sgemm_tcu-5:simx": 89, + "tensor:sgemm_tcu-6:simx": 88, + "tensor:sgemm_tcu-7:simx": 88, + "tensor:sgemm_tcu-8:rtlsim": 45, + "tensor:sgemm_tcu-9:rtlsim": 27, + "tensor:sgemm_tcu-fedp2k:simx": 89, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt16:rtlsim": 73, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt32:rtlsim": 159, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt4:rtlsim": 38, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxbf8-nt8:rtlsim": 43, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt16:rtlsim": 87, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt32:rtlsim": 172, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt4:rtlsim": 41, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp4-nt8:rtlsim": 47, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt16:rtlsim": 70, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt32:rtlsim": 147, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt4:rtlsim": 38, + "tensor_mx:sgemm_tcu_mx-rtlsim-mxfp8-nt8:rtlsim": 45, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt16:rtlsim": 86, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt32:rtlsim": 189, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt4-tls:rtlsim": 43, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt4:rtlsim": 39, + "tensor_mx:sgemm_tcu_mx-rtlsim-nvfp4-nt8:rtlsim": 45, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt16:simx": 74, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt32:simx": 72, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt4:simx": 75, + "tensor_mx:sgemm_tcu_mx-simx-mxbf8-nt8:simx": 76, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt16:simx": 72, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt32:simx": 74, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt4:simx": 73, + "tensor_mx:sgemm_tcu_mx-simx-mxfp4-nt8:simx": 73, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt16:simx": 74, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt32:simx": 74, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt4:simx": 87, + "tensor_mx:sgemm_tcu_mx-simx-mxfp8-nt8:simx": 75, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt16:simx": 72, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt32:simx": 72, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt4-tls:simx": 73, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt4:simx": 76, + "tensor_mx:sgemm_tcu_mx-simx-nvfp4-nt8:simx": 73, + "tensor_sp:sgemm_tcu_sp-10:simx": 90, + "tensor_sp:sgemm_tcu_sp-11:simx": 90, + "tensor_sp:sgemm_tcu_sp-12:simx": 90, + "tensor_sp:sgemm_tcu_sp-13:rtlsim": 33, + "tensor_sp:sgemm_tcu_sp-14:rtlsim": 33, + "tensor_sp:sgemm_tcu_sp-15:rtlsim": 34, + "tensor_sp:sgemm_tcu_sp-16:rtlsim": 35, + "tensor_sp:sgemm_tcu_sp-17:simx": 90, + "tensor_sp:sgemm_tcu_sp-18:simx": 90, + "tensor_sp:sgemm_tcu_sp-19:simx": 90, + "tensor_sp:sgemm_tcu_sp-1:simx": 98, + "tensor_sp:sgemm_tcu_sp-20:simx": 91, + "tensor_sp:sgemm_tcu_sp-21:rtlsim": 38, + "tensor_sp:sgemm_tcu_sp-22:rtlsim": 39, + "tensor_sp:sgemm_tcu_sp-23:rtlsim": 39, + "tensor_sp:sgemm_tcu_sp-24:rtlsim": 39, + "tensor_sp:sgemm_tcu_sp-25:simx": 90, + "tensor_sp:sgemm_tcu_sp-26:simx": 91, + "tensor_sp:sgemm_tcu_sp-27:simx": 89, + "tensor_sp:sgemm_tcu_sp-28:simx": 89, + "tensor_sp:sgemm_tcu_sp-29:rtlsim": 75, + "tensor_sp:sgemm_tcu_sp-2:simx": 89, + "tensor_sp:sgemm_tcu_sp-30:rtlsim": 63, + "tensor_sp:sgemm_tcu_sp-31:rtlsim": 73, + "tensor_sp:sgemm_tcu_sp-32:rtlsim": 64, + "tensor_sp:sgemm_tcu_sp-33:simx": 89, + "tensor_sp:sgemm_tcu_sp-34:simx": 90, + "tensor_sp:sgemm_tcu_sp-35:simx": 90, + "tensor_sp:sgemm_tcu_sp-36:simx": 90, + "tensor_sp:sgemm_tcu_sp-37:rtlsim": 129, + "tensor_sp:sgemm_tcu_sp-38:rtlsim": 153, + "tensor_sp:sgemm_tcu_sp-39:rtlsim": 155, + "tensor_sp:sgemm_tcu_sp-3:simx": 91, + "tensor_sp:sgemm_tcu_sp-40:rtlsim": 139, + "tensor_sp:sgemm_tcu_sp-4:simx": 91, + "tensor_sp:sgemm_tcu_sp-5:rtlsim": 53, + "tensor_sp:sgemm_tcu_sp-6:rtlsim": 27, + "tensor_sp:sgemm_tcu_sp-7:rtlsim": 28, + "tensor_sp:sgemm_tcu_sp-8:rtlsim": 28, + "tensor_sp:sgemm_tcu_sp-9:simx": 90, + "tensor_sp_mx:sgemm_tcu_sp_mx-rtlsim-mxfp8-nt4:rtlsim": 27, + "tensor_sp_mx:sgemm_tcu_sp_mx-rtlsim-mxfp8-nt8:rtlsim": 27, + "tensor_sp_mx:sgemm_tcu_sp_mx-simx-mxfp8-nt4:simx": 93, + "tensor_sp_mx:sgemm_tcu_sp_mx-simx-mxfp8-nt8:simx": 91, + "tensor_wg:sgemm_tcu_wg-10:simx": 91, + "tensor_wg:sgemm_tcu_wg-1:simx": 92, + "tensor_wg:sgemm_tcu_wg-2:simx": 91, + "tensor_wg:sgemm_tcu_wg-3:simx": 90, + "tensor_wg:sgemm_tcu_wg-4:simx": 91, + "tensor_wg:sgemm_tcu_wg-5:rtlsim": 149, + "tensor_wg:sgemm_tcu_wg-6:rtlsim": 1694, + "tensor_wg:sgemm_tcu_wg-7:rtlsim": 4182, + "tensor_wg:sgemm_tcu_wg-8:simx": 91, + "tensor_wg:sgemm_tcu_wg-9:rtlsim": 143, + "tensor_wg:sgemm_tcu_wg-fedp2k:simx": 90, + "tensor_wg:sgemm_tcu_wg_mx-baseline:simx": 95, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-nvfp4:simx": 92, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-rs:simx": 91, + "tensor_wg:sgemm_tcu_wg_mx-fedp2k-ss:simx": 92, + "tensor_wg:sgemm_tcu_wg_sp-1:simx": 93, + "tensor_wg:sgemm_tcu_wg_sp-2:simx": 93, + "tensor_wg:sgemm_tcu_wg_sp-3:simx": 93, + "tensor_wg:sgemm_tcu_wg_sp-4:rtlsim": 156, + "tensor_wg:sgemm_tcu_wg_sp-5:rtlsim": 1688, + "tensor_wg:sgemm_tcu_wg_sp-6:rtlsim": 3459, + "tensor_wg:sgemm_tcu_wg_sp-7:simx": 92, + "tensor_wg:sgemm_tcu_wg_sp-8:rtlsim": 158, + "tensor_wg:sgemm_tcu_wg_sp-fedp2k-rs:simx": 94, + "tensor_wg:sgemm_tcu_wg_sp-fedp2k-ss:simx": 93, + "unittest:hw-fpu:None": 2, + "unittest:hw-tcu-dsp:None": 71, + "unittest:hw-tcu:None": 227, + "unittest:hw-tex-lerp:None": 1, + "unittest:hw:None": 256, + "unittest:sw:None": 176, + "vm:isa-10:rtlsim": 45, + "vm:isa-10:simx": 113, + "vm:isa-1:rtlsim": 68, + "vm:isa-1:simx": 193, + "vm:isa-2:rtlsim": 6, + "vm:isa-2:simx": 3, + "vm:isa-3:rtlsim": 4, + "vm:isa-3:simx": 3, + "vm:isa-4:rtlsim": 16, + "vm:isa-4:simx": 9, + "vm:isa-5:rtlsim": 44, + "vm:isa-5:simx": 112, + "vm:isa-6:rtlsim": 65, + "vm:isa-6:simx": 188, + "vm:isa-7:rtlsim": 6, + "vm:isa-7:simx": 3, + "vm:isa-8:rtlsim": 4, + "vm:isa-8:simx": 3, + "vm:isa-9:rtlsim": 15, + "vm:isa-9:simx": 9, + "vm:mc-l2-amo:rtlsim": 955, + "vm:mc-l2-amo:simx": 194, + "vm:mc-l3:rtlsim": 429, + "vm:mc-l3:simx": 196, + "vm:tlb-concurrent:rtlsim": 201, + "vm:tlb-concurrent:simx": 188, + "vm:tlb-stress-l1:rtlsim": 28, + "vm:tlb-stress-l1:simx": 94, + "vm:tlb-stress-l2:rtlsim": 0, + "vm:tlb-stress-l2:simx": 0, + "vm:tlb-stress-walk:rtlsim": 1, + "vm:tlb-stress-walk:simx": 1, + "vm:vm-amo:rtlsim": 58, + "vm:vm-amo:simx": 189, + "vm:vm-drain:rtlsim": 0, + "vm:vm-drain:simx": 0, + "vm:vm-fault:rtlsim": 28, + "vm:vm-fault:simx": 99, + "vm:vm-fence:rtlsim": 0, + "vm:vm-fence:simx": 0, + "vm:vm-stress:rtlsim": 1, + "vm:vm-stress:simx": 1, + "vm:vm-superpage:rtlsim": 28, + "vm:vm-superpage:simx": 95, + "vulkan:ff-all-sw:simx": 93, + "vulkan:ff-om-only:simx": 102, + "vulkan:ff-raster-om:simx": 108, + "vulkan:ff-raster-only:simx": 100, + "vulkan:ff-raster-tex-om:simx": 111, + "vulkan:ff-raster-tex:simx": 105, + "vulkan:ff-tex-om:simx": 114, + "vulkan:ff-tex-only:simx": 102, + "vulkan:gfx-multidraw:rtlsim": 40, + "vulkan:gfx-multidraw:simx": 108, + "vulkan:gfx-rt-rtquery-id:simx": 126, + "vulkan:isa-1:simx": 1756, + "vulkan:isa-4:xrtsim": 1684, + "vulkan:rt-raytrace:rtlsim": 96, + "vulkan:rt-raytrace:simx": 123, + "vulkan:rt-rtquery:simx": 123 + }, + "source": { + "run": 37241939097, + "sha": "8524ae69c8237c7c1878742171e1b13968622acf" + } +} diff --git a/ci/testcase.py b/ci/testcase.py index 7bd506c109..0e8da16c65 100644 --- a/ci/testcase.py +++ b/ci/testcase.py @@ -17,12 +17,15 @@ JSON (category x driver x xlen) cells for the GitHub matrix testcase.py select --changed-from=REF categories whose touches[] intersect the diff (path-scaling) + testcase.py expr --cell=CATEGORY --driver=DRIVER + the pytest -m expression a matrix cell runs """ import argparse import copy import glob import json +import math import os import subprocess import sys @@ -37,6 +40,7 @@ # Execution driver name -> its sim source directory under sim/. _DRIVER_TO_SIMDIR = {"simx": "simx", "rtlsim": "rtlsim", "xrt": "xrtsim", "opae": "opaesim"} VALID_DRIVERS = set(_DRIVER_TO_SIMDIR) +DRIVER_MARKERS = sorted({_DRIVER_TO_MARKER.get(d, d) for d in VALID_DRIVERS}) VALID_VIA = {"blackbox", "make-run", "script"} VALID_CHECK = {"model_parity", "perf_gate"} @@ -51,6 +55,15 @@ # simx<->rtlsim cycle-parity default: both timing models must agree within 5%. DEFAULT_PARITY_TOLERANCE = 0.05 +# Per-case wall times (seconds, per xlen) measured by a CI run, recorded by +# ci/record_runtimes.py from the run's junit reports. A cell whose recorded total +# exceeds SHARD_BUDGET_S is split into shards balanced on them, so no cell runs +# into the tests job's timeout. +RUNTIMES_FILE = os.path.join(os.path.dirname(os.path.abspath(__file__)), "runtimes.json") +SHARD_BUDGET_S = 150 * 60 +# Weight of a case with no recorded time (new, or not yet run on that xlen). +DEFAULT_CASE_S = 120 + def driver_marker(driver): """User-facing marker name for an execution driver (xrt->xrtsim, ...).""" @@ -128,6 +141,19 @@ def build_key(self): """ return (self.driver, self.configs) + def sim_signature(self, xlen): + """What a run of this case compiles its sim with, or None if it builds + none. Wider than build_key(): a blackbox run rebuilds the sim from its + full flag set (configs plus the shape and extra flags), so two cases with + equal configs but different shapes still rebuild between them. + """ + if not self.needs_sim: + return None + configs = _subst(self.configs, xlen) + if self.via == "blackbox": + return (self.driver, configs, tuple(_shape_flags(self.shape)), self.flags) + return (self.driver, configs) + def markers(self): """pytest marker names for `-m` selection (one per value).""" m = [self.category, self.tier] @@ -370,6 +396,61 @@ def _filter(cases, args): return out +def load_runtimes(xlen): + try: + with open(RUNTIMES_FILE) as fh: + return json.load(fh).get(str(xlen), {}) + except OSError: + return {} + + +def cell_name(case): + """The matrix cell a case runs in: its check's cell for a check case (see + cmd_matrix), its category's otherwise.""" + return case.check or case.category + + +def shard_plan(case_ids, xlen, shards=None): + """Split a cell's cases into shards balanced on their recorded runtimes. + + Longest case first onto the least-loaded shard, with ties broken on case id + and shard index, so the planner and every shard derive the same split from + the same ids. Unless `shards` is given, the count is the fewest that keep + the cell's recorded total within SHARD_BUDGET_S per shard. + Returns (shards, {case id: 0-based shard}). + """ + runtimes = load_runtimes(xlen) + weight = {i: runtimes.get(i, DEFAULT_CASE_S) for i in case_ids} + if shards is None: + shards = max(1, math.ceil(sum(weight.values()) / SHARD_BUDGET_S)) + load = [0.0] * shards + plan = {} + for i in sorted(weight, key=lambda i: (-weight[i], i)): + k = min(range(shards), key=lambda s: (load[s], s)) + plan[i] = k + load[k] += weight[i] + return shards, plan + + +def cell_expr(name, driver): + """The pytest -m expression a matrix cell runs. + + A host cell holds a category's driverless cases only: its driver cases run in + their own driver's cell, and without the exclusion a host cell re-ran all of + them. A check's cases run once, in the check's cell, so every other cell + excludes them. + """ + expr = name + if driver == "host": + expr += " and not ({})".format(" or ".join(DRIVER_MARKERS)) + else: + expr += " and " + driver + for check in sorted(VALID_CHECK): + if check != name: + expr += " and not " + check + return expr + + def cmd_matrix(args): # One GitHub matrix cell per (category, driver, xlen): the build tree is # per-xlen, so xlen is flattened out (not a per-cell list). @@ -379,28 +460,50 @@ def cmd_matrix(args): # from its own category's cell so it does not run twice. That cell is keyed on # the CHECK, not on a category that happens to share its name -- a check is a # marker across suites, and no category name may control whether it runs. + # + # A cell whose recorded runtime exceeds SHARD_BUDGET_S is emitted as `shards` + # cells, each running shard `shard` of the split (shard_plan). The split is + # over every case the cell's -m expression selects, not only those this + # event's filters planned the cell for, because that is what the cell runs. xfilter = {int(x) for x in args.xlen.split(",")} if getattr(args, "xlen", None) else None + cases = load_all() cells = {} - for c in _filter(load_all(), args): + for c in _filter(cases, args): drv = c.marker_driver or "host" for xlen in c.xlens: if xfilter and xlen not in xfilter: continue - name = c.check or c.category + name = cell_name(c) key = (name, drv, xlen) cell = cells.setdefault(key, { "category": name, "driver": drv, "xlen": xlen, "needs": set(), }) cell["needs"].update(c.needs) + members = {} + for c in cases: + if c.tier in OPT_IN_TIERS: + continue + for xlen in c.xlens: + members.setdefault((cell_name(c), c.marker_driver or "host", xlen), []).append(c.id) out = [] - for cell in cells.values(): + for key, cell in cells.items(): cell["needs"] = sorted(cell["needs"]) - out.append(cell) - out.sort(key=lambda c: (c["category"], c["driver"], c["xlen"])) + shards, _ = shard_plan(members.get(key, []), key[2]) + if shards == 1: + out.append(cell) + continue + for k in range(shards): + out.append(dict(cell, shard=k + 1, shards=shards)) + out.sort(key=lambda c: (c["category"], c["driver"], c["xlen"], c.get("shard", 0))) print(json.dumps(out)) return 0 +def cmd_expr(args): + print(cell_expr(args.cell, args.driver)) + return 0 + + def cmd_select(args): print(" ".join(sorted({c.category for c in _filter(load_all(), args)}))) return 0 @@ -515,6 +618,11 @@ def main(argv=None): sub.add_parser("checks", help="cross-cutting check names (one cell each)" ).set_defaults(func=cmd_checks) + e = sub.add_parser("expr", help="the pytest -m expression a matrix cell runs") + e.add_argument("--cell", required=True, help="category or check name") + e.add_argument("--driver", required=True, help="driver marker, or host") + e.set_defaults(func=cmd_expr) + sub.add_parser("lint", help="validate the test cases").set_defaults(func=cmd_lint) args = p.parse_args(argv) From ba32c46627865a4bb4618ea663ebad0ad6d04c27 Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 11:14:10 -0700 Subject: [PATCH 6/8] ci: hold asic_gate Fmax to the target clock, not to the baseline The Yosys gate now gates only cell area against the baseline; Fmax is held to the build's clock target by the existing BELOW TARGET check. ABC maps against that target (-D) and spends the slack beyond it recovering area, so above the target Fmax moves with whatever else the design contains. After #424 the RTU's worst path is an FMA stage (tri_pe fsub_r, align -> add -> LZC) in VX_fma_unit_rtl, which #424 did not touch, yet it moved 957 -> 903 MHz against a 400 MHz target and failed the baseline check. The fpga_gate is unchanged: Vivado reports what the routed design achieves. Co-Authored-By: Claude Opus 5.5 --- ci/synth_gate.py | 7 ++++++- ci/testcases/asic_gate.yaml | 7 ++++--- docs/designs/continuous_integration.md | 10 ++++++---- 3 files changed, 16 insertions(+), 8 deletions(-) diff --git a/ci/synth_gate.py b/ci/synth_gate.py index f9e1dec96c..820089a69d 100755 --- a/ci/synth_gate.py +++ b/ci/synth_gate.py @@ -303,7 +303,12 @@ class Yosys(Tool): metrics = ("fmax_mhz", "wns_ns", "tns_ns", "cell_area_um2", "seq_area_um2", "sram_area_um2", "cell_count", "power_mw", "build_time_s") - gated = ("fmax_mhz", "cell_area_um2") + # Fmax is held to the build's clock target (gate()'s BELOW TARGET check), + # not to its baseline. ABC maps against that target (-D) and spends any + # slack beyond it recovering area, so above the target Fmax moves with + # whatever else the design contains: #424 left the RTU's worst path in an + # unchanged FMA stage yet moved it 957 -> 903 MHz. + gated = ("cell_area_um2",) integral = ("cell_count", "build_time_s") columns = (("fmax", "fmax_mhz", "%.0f", 9), ("area_um2", "cell_area_um2", "%.0f", 11), diff --git a/ci/testcases/asic_gate.yaml b/ci/testcases/asic_gate.yaml index 8a8c0dd34a..5f81318062 100644 --- a/ci/testcases/asic_gate.yaml +++ b/ci/testcases/asic_gate.yaml @@ -1,6 +1,7 @@ -# ASIC synthesis-regression gate: post-synthesis Fmax/cell area must not regress -# past a threshold vs. a checked-in golden result -# (ci/baselines/synthesis/yosys/.json). +# ASIC synthesis-regression gate: post-synthesis cell area must not regress past +# a threshold vs. a checked-in golden result +# (ci/baselines/synthesis/yosys/.json), and Fmax must meet the build's +# target clock. # # Sibling of the fpga_gate (fpga_gate.yaml): same discipline — golden baseline, # hard threshold, ratcheted improvements, results only ever re-recorded by a diff --git a/docs/designs/continuous_integration.md b/docs/designs/continuous_integration.md index 04238ea748..c959934e45 100644 --- a/docs/designs/continuous_integration.md +++ b/docs/designs/continuous_integration.md @@ -260,7 +260,7 @@ and assert the results against golden baselines. |---|---|---| | flow | Vivado synthesis, place and route | sv2v → Yosys → ABC → OpenSTA, on ASAP7 | | design tree | `hw/syn/xilinx/dut` | `hw/syn/yosys/dut` | -| gated | `fmax_mhz`, `lut` | `fmax_mhz`, `cell_area_um2` | +| gated against the baseline | `fmax_mhz`, `lut` | `cell_area_um2` (Fmax only against the target clock) | | also recorded | `wns_ns`, `lutram`, `ff`, `bram`, `uram`, `dsp`, critical paths, high-fanout nets | `wns_ns`, `tns_ns`, `seq_area_um2`, `sram_area_um2`, `cell_count`, `power_mw` | | comparable when | same device, optimization level, xlen | same PDK, VT, corner, xlen | | runner | self-hosted, licensed Vivado | hosted, one job per design | @@ -318,9 +318,11 @@ what the implemented design achieved. ABC maps to the target period and stops, so a Yosys design that closes does so with picoseconds of margin and its Fmax sits just above its clock by construction. On the ASIC side **cell area is the sensitive metric** and Fmax is mostly met-or-missed — -which is what the target-clock assertion checks. It also makes the slack -source matter: worst *negative* slack clamps at zero, so the flow reports the -signed worst slack instead. +which is what the target-clock assertion checks. Above the target, ABC spends +the slack recovering area, so where Fmax lands moves with everything else in the +design; the ASIC gate therefore holds Fmax to the target and not to a baseline. +It also makes the slack source matter: worst *negative* slack clamps at zero, +so the flow reports the signed worst slack instead. | verdict | meaning | fails the run | |---|---|---| From f1eb8601cedc2517baa8c33db705f22cde40c199 Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 05:50:23 -0700 Subject: [PATCH 7/8] ci: re-record the rtu asic_gate baseline after Vulkan rt (#424) #424 grew the RTU by 8.0% cell area (91600 -> 98965 um^2) and SRAM area by 10%: the per-context state its Vulkan ray-tracing features add (candidate records, object-ray staging, the any-hit/intersection resume state) and the watertight triangle PE. Measured by the nightly on 3c3a84b40 (run 37263315588, Yosys 0.69 / OpenSTA 3.1.0, config_hash unchanged) and recorded from that run's report; the RTU RTL is the same on this branch. Fmax 903 MHz against the 400 MHz target. Co-Authored-By: Claude Opus 5.5 --- ci/baselines/synthesis/yosys/graphics.json | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/ci/baselines/synthesis/yosys/graphics.json b/ci/baselines/synthesis/yosys/graphics.json index e68752040c..31b1951e13 100644 --- a/ci/baselines/synthesis/yosys/graphics.json +++ b/ci/baselines/synthesis/yosys/graphics.json @@ -83,9 +83,9 @@ "configs": "", "dut": "rtu", "result": { - "build_time_s": 880, - "cell_area_um2": 91600.089, - "cell_count": 842051, + "build_time_s": 1001, + "cell_area_um2": 98964.783, + "cell_count": 928942, "env": { "corner": "tt", "pdk": "asap7", @@ -95,12 +95,12 @@ "xlen": 32, "yosys": "Yosys 0.69" }, - "fmax_mhz": 956.755, - "power_mw": 349.0, - "seq_area_um2": 36179.978, - "sram_area_um2": 10949.6, + "fmax_mhz": 903.424, + "power_mw": 361.0, + "seq_area_um2": 37464.185, + "sram_area_um2": 12099.2, "tns_ns": 0.0, - "wns_ns": 1.455 + "wns_ns": 1.393 } }, "tex": { From 62aaa310ea32341d63f8f63012f454d53fd80b74 Mon Sep 17 00:00:00 2001 From: tinebp Date: Mon, 5 Oct 2026 08:49:26 -0700 Subject: [PATCH 8/8] ci: re-record the fpga_gate top and gfx baselines Both DUTs contain RTL this branch changes: top the CP AXI-Lite regfile (flat decode), gfx the raster block evaluator (early-Z plane register). Recorded with Vivado 2024.2, config_hash unchanged; the previous baselines also predate Vulkan rt (#424), which reshaped the RTU's memories. top 250 MHz target: 252.6 -> 250.9 MHz, LUT 136066 -> 136762 (+0.5%) gfx 250 MHz target: 250.1 -> 250.1 MHz, LUT 202801 -> 204871 (+1.0%), FF +3%, LUTRAM 10287 -> 16721, BRAM 312 -> 165 (RTU memories now map to LUTRAM) Co-Authored-By: Claude Opus 5.5 --- ci/baselines/synthesis/xilinx/core.json | 176 +++++++++--------- ci/baselines/synthesis/xilinx/graphics.json | 190 ++++++++++---------- 2 files changed, 183 insertions(+), 183 deletions(-) diff --git a/ci/baselines/synthesis/xilinx/core.json b/ci/baselines/synthesis/xilinx/core.json index 1c78fe4ad7..70cda5d27a 100644 --- a/ci/baselines/synthesis/xilinx/core.json +++ b/ci/baselines/synthesis/xilinx/core.json @@ -360,107 +360,107 @@ "dut": "top", "result": { "bram": 80, - "build_time_s": 3513, + "build_time_s": 3788, "critical_paths": [ { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/pe_switch/rsp_arb/g_input_select.g_arbiter.g_out_buf[0].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[24]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/dcache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[0].bank/core_rsp_queue/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[28]/D", "group": "core_clock", - "levels": 1, - "logic_ns": 0.115, - "route_ns": 3.825, - "route_pct": 97.081, - "slack_ns": 0.041, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/csr_unit/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/C" + "levels": 10, + "logic_ns": 0.912, + "route_ns": 3.054, + "route_pct": 77.006, + "slack_ns": 0.014, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/dcache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[0].bank/g_amo.amo/g_commit.wbq_addr_reg[3][21]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/dmmu/g_pipe[1].out_pipe/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[119]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[35]/CE", "group": "core_clock", - "levels": 9, - "logic_ns": 0.965, - "route_ns": 2.929, - "route_pct": 75.22, - "slack_ns": 0.085, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[101]/C" + "levels": 10, + "logic_ns": 1.145, + "route_ns": 2.714, + "route_pct": 70.332, + "slack_ns": 0.037, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/dcache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[1].bank/core_rsp_queue/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[43]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[88]/CE", "group": "core_clock", "levels": 10, - "logic_ns": 0.923, - "route_ns": 2.971, - "route_pct": 76.295, - "slack_ns": 0.087, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/dcache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[1].bank/reg_cmt/g_pipe.g_partial_reset.pipe_free_reg[0][73]/C" + "logic_ns": 1.145, + "route_ns": 2.714, + "route_pct": 70.332, + "slack_ns": 0.037, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/pe_switch/rsp_arb/g_input_select.g_arbiter.g_out_buf[0].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.buffer_r_reg[55]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[35]/CE", "group": "core_clock", - "levels": 1, - "logic_ns": 0.167, - "route_ns": 3.726, - "route_pct": 95.71, - "slack_ns": 0.09, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/csr_unit/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/C" + "levels": 10, + "logic_ns": 1.202, + "route_ns": 2.651, + "route_pct": 68.802, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[105]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/pe_switch/rsp_arb/g_input_select.g_arbiter.g_out_buf[0].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[103]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[88]/CE", "group": "core_clock", - "levels": 1, - "logic_ns": 0.231, - "route_ns": 3.658, - "route_pct": 94.06, - "slack_ns": 0.093, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/csr_unit/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/C" + "levels": 10, + "logic_ns": 1.202, + "route_ns": 2.651, + "route_pct": 68.802, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[105]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/pe_switch/rsp_arb/g_input_select.g_arbiter.g_out_buf[0].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[28]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[1]/CE", "group": "core_clock", - "levels": 1, - "logic_ns": 0.114, - "route_ns": 3.772, - "route_pct": 97.067, - "slack_ns": 0.095, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/csr_unit/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/C" + "levels": 10, + "logic_ns": 1.145, + "route_ns": 2.709, + "route_pct": 70.293, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/pe_switch/rsp_arb/g_input_select.g_arbiter.g_out_buf[0].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[87]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[2]/CE", "group": "core_clock", - "levels": 1, - "logic_ns": 0.179, - "route_ns": 3.705, - "route_pct": 95.391, - "slack_ns": 0.099, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/sfu_unit/csr_unit/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/C" + "levels": 10, + "logic_ns": 1.145, + "route_ns": 2.709, + "route_pct": 70.293, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/dmmu/g_pipe[1].out_pipe/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[113]/D", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[54]/CE", "group": "core_clock", - "levels": 9, - "logic_ns": 1.008, - "route_ns": 2.869, - "route_pct": 74.0, - "slack_ns": 0.104, - "startpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[107]/C" + "levels": 10, + "logic_ns": 1.145, + "route_ns": 2.709, + "route_pct": 70.293, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/icache/g_cache_wrap[0].cache_wrap/g_cache.cache/mem_req_xbar/g_fallback.xbar_switch/g_passthru.out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg/S", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[94]/CE", "group": "core_clock", - "levels": 0, - "logic_ns": 0.077, - "route_ns": 3.699, - "route_pct": 97.961, - "slack_ns": 0.106, - "startpoint": "vx_reset_shift_r_reg[7]_replica_15/C" + "levels": 10, + "logic_ns": 1.145, + "route_ns": 2.709, + "route_pct": 70.293, + "slack_ns": 0.042, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[104]/C" }, { - "endpoint": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/icache/g_cache_wrap[0].cache_wrap/g_cache.cache/mem_req_xbar/g_fallback.xbar_switch/g_passthru.out_buf/g_eb2.stream_buffer/g_buffer.valid_out_r_reg/R", + "endpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[1]/CE", "group": "core_clock", - "levels": 0, - "logic_ns": 0.077, - "route_ns": 3.696, - "route_pct": 97.959, - "slack_ns": 0.109, - "startpoint": "vx_reset_shift_r_reg[7]_replica_15/C" + "levels": 10, + "logic_ns": 1.202, + "route_ns": 2.646, + "route_pct": 68.761, + "slack_ns": 0.047, + "startpoint": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/mem_unit/g_dcache_adapters[0].dcache_adapter/stream_unpack/g_unpack.g_outbuf[1].out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[105]/C" } ], "dsp": 48, @@ -470,64 +470,64 @@ "vivado": "2024.2", "xlen": 32 }, - "ff": 143375, - "fmax_mhz": 252.6, + "ff": 145017, + "fmax_mhz": 250.9, "high_fanout_nets": [ { "driver": "LUT2", "fanout": 5638, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.div_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.div_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "BUFGCE", "fanout": 5633, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.div_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.div_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "LUT2", "fanout": 4717, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fma.pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fma.pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "BUFGCE", "fanout": 4669, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fma.pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fma.pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "LUT2", "fanout": 3318, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.sqrt_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[1].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.sqrt_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "BUFGCE", "fanout": 3301, - "net": "vortex_axi/vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.sqrt_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + "net": "g_multi_port.vortex/g_clusters[0].cluster/g_sockets[0].socket/g_cores[0].core/execute/fpu_unit/g_blocks[0].fpu_dsp/g_fdivsqrt.sqrt_pe_ser/out_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" + }, + { + "driver": "BUFGCE", + "fanout": 2003, + "net": "cp_core/u_regfile/reset_0[0]" }, { "driver": "MUXF7", "fanout": 1705, - "net": "vortex_axi/vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_2" + "net": "g_multi_port.vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_2" }, { "driver": "MUXF7", "fanout": 1682, - "net": "vortex_axi/vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_5" + "net": "g_multi_port.vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_5" }, { "driver": "MUXF7", "fanout": 1537, - "net": "vortex_axi/vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_3" - }, - { - "driver": "MUXF7", - "fanout": 1476, - "net": "vortex_axi/vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_4" + "net": "g_multi_port.vortex/g_clusters[0].cluster/l2tlb/mshr/g_buffer.g_no_out_reg.data_out_r_reg[32]_3" } ], - "lut": 136066, + "lut": 136762, "lutram": 9536, "uram": 0, - "wns_ns": 0.041 + "wns_ns": 0.014 } } } diff --git a/ci/baselines/synthesis/xilinx/graphics.json b/ci/baselines/synthesis/xilinx/graphics.json index 1d1d44464d..af42d515fb 100644 --- a/ci/baselines/synthesis/xilinx/graphics.json +++ b/ci/baselines/synthesis/xilinx/graphics.json @@ -5,129 +5,129 @@ "configs": "-DVX_CFG_NUM_THREADS=16 -DVX_CFG_NUM_WARPS=16", "dut": "gfx", "result": { - "bram": 312, - "build_time_s": 5285, + "bram": 165, + "build_time_s": 6961, "critical_paths": [ { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[3][0][14]/CE", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_8/DINADIN[25]", "group": "core_clock", "levels": 8, - "logic_ns": 0.964, - "route_ns": 2.93, - "route_pct": 75.241, + "logic_ns": 0.774, + "route_ns": 2.922, + "route_pct": 79.061, "slack_ns": 0.002, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[1]/C" + "startpoint": "g_rtu_core[0].rtu_core/scheduler/structaddr_q_reg[5]_rep__6/C" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[3][0][14]/CE", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_9/DINADIN[31]", "group": "core_clock", - "levels": 9, - "logic_ns": 1.063, - "route_ns": 2.82, - "route_pct": 72.623, - "slack_ns": 0.012, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[4]/C" + "levels": 8, + "logic_ns": 0.916, + "route_ns": 2.767, + "route_pct": 75.132, + "slack_ns": 0.003, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/structaddr_q_reg[5]_rep__6/C" }, { - "endpoint": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/g_rsp_N.g_rsp_full.g_rsp_store[0].g_j[54].rsp_store_reg_0_7_16_16/SP/I", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_10/DINBDIN[16]", "group": "core_clock", - "levels": 0, - "logic_ns": 0.081, - "route_ns": 3.827, - "route_pct": 97.927, - "slack_ns": 0.015, - "startpoint": "tcache/g_cache_wrap[0].cache_wrap/g_cache.cache/core_rsp_xbar/g_fallback.xbar_switch/g_passthru.out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[28]/C" + "levels": 8, + "logic_ns": 0.921, + "route_ns": 2.767, + "route_pct": 75.029, + "slack_ns": 0.003, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/structaddr_q_reg[5]_rep__6/C" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[3][0][14]/CE", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_10/DINBDIN[7]", "group": "core_clock", - "levels": 8, - "logic_ns": 0.943, - "route_ns": 2.936, - "route_pct": 75.689, - "slack_ns": 0.016, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[10]/C" + "levels": 13, + "logic_ns": 1.296, + "route_ns": 2.395, + "route_pct": 64.888, + "slack_ns": 0.004, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/rewalk_q_reg/C" }, { - "endpoint": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/g_rsp_N.g_rsp_full.g_rsp_store[0].g_j[47].rsp_store_reg_0_7_29_29/SP/I", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_arbiter/g_model1.reqs_mask_reg[15]/CE", "group": "core_clock", - "levels": 0, - "logic_ns": 0.08, - "route_ns": 3.827, - "route_pct": 97.952, - "slack_ns": 0.016, - "startpoint": "tcache/g_cache_wrap[0].cache_wrap/g_cache.cache/core_rsp_xbar/g_fallback.xbar_switch/g_passthru.out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[41]/C" + "levels": 14, + "logic_ns": 1.569, + "route_ns": 2.322, + "route_pct": 59.675, + "slack_ns": 0.004, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/wake_self_r_reg/C" }, { - "endpoint": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/rsp_buf/g_eb1.pipe_buffer/g_register.g_pipe_regs[0].pipe_register/g_pipe.g_partial_reset.pipe_reg[0][1842]/D", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_10/DINBDIN[3]", "group": "core_clock", - "levels": 1, - "logic_ns": 0.17, - "route_ns": 3.794, - "route_pct": 95.712, - "slack_ns": 0.018, - "startpoint": "tcache/g_cache_wrap[0].cache_wrap/g_cache.cache/core_rsp_xbar/g_fallback.xbar_switch/g_passthru.out_buf/g_eb2.stream_buffer/g_buffer.g_out_reg.data_out_r_reg[42]/C" + "levels": 10, + "logic_ns": 0.813, + "route_ns": 2.896, + "route_pct": 78.082, + "slack_ns": 0.004, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/rewalk_q_reg/C" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordoff_reg[9][2][19]/D", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/bw_job_done_reg_rep__0/D", "group": "core_clock", "levels": 9, - "logic_ns": 0.895, - "route_ns": 3.067, - "route_pct": 77.41, - "slack_ns": 0.019, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[10]/C" + "logic_ns": 1.71, + "route_ns": 2.218, + "route_pct": 56.471, + "slack_ns": 0.005, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/window_store/g_no_asic.g_sync.g_bram.g_write_first.g_wren.ram_reg_2/CLKARDCLK" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[3][0][14]/CE", + "endpoint": "graphics/ocache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[0].bank/cache_data/g_data_slice[3].data_store/g_no_asic.g_sync.g_bram.g_read_first.g_wren.ram_reg/ENBWREN", "group": "core_clock", - "levels": 9, - "logic_ns": 1.006, - "route_ns": 2.864, - "route_pct": 74.006, - "slack_ns": 0.026, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[6]/C" + "levels": 11, + "logic_ns": 0.909, + "route_ns": 2.718, + "route_pct": 74.939, + "slack_ns": 0.005, + "startpoint": "graphics/ocache/g_cache_wrap[0].cache_wrap/g_cache.cache/g_banks[0].bank/reg_s0/g_pipe.g_partial_reset.pipe_free_reg[0][78]_replica/C" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[15][3][6]/D", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/ctx_store/g_no_asic.g_sync.g_auto.g_write_first.g_no_wren.ram_reg_8/DINBDIN[17]", "group": "core_clock", - "levels": 9, - "logic_ns": 0.832, - "route_ns": 3.123, - "route_pct": 78.962, - "slack_ns": 0.026, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[10]/C" + "levels": 8, + "logic_ns": 0.788, + "route_ns": 2.948, + "route_pct": 78.907, + "slack_ns": 0.006, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/structaddr_q_reg[5]_rep__6/C" }, { - "endpoint": "g_rtu_core[0].rtu_core/scheduler/coll_ordt_reg[3][0][14]/CE", + "endpoint": "g_rtu_core[0].rtu_core/scheduler/bw_job_done_reg_rep__0/D", "group": "core_clock", - "levels": 8, - "logic_ns": 1.038, - "route_ns": 2.828, - "route_pct": 73.148, - "slack_ns": 0.029, - "startpoint": "g_rtu_core[0].rtu_core/scheduler/box_coll_hot_reg[4]/C" + "levels": 10, + "logic_ns": 1.804, + "route_ns": 2.123, + "route_pct": 54.066, + "slack_ns": 0.006, + "startpoint": "g_rtu_core[0].rtu_core/scheduler/window_store/g_no_asic.g_sync.g_bram.g_write_first.g_wren.ram_reg_2/CLKARDCLK" } ], - "dsp": 758, + "dsp": 752, "env": { "device": "xcu55c-fsvh2892-2L-e", "opt_level": 3, "vivado": "2024.2", "xlen": 32 }, - "ff": 214971, + "ff": 221979, "fmax_mhz": 250.1, "high_fanout_nets": [ { - "driver": "LUT2", - "fanout": 7655, - "net": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/req_ibuf/allocator/req_ready_bufg_place" + "driver": "BUFGCE", + "fanout": 9725, + "net": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/req_ibuf/allocator/req_ready" }, { "driver": "BUFGCE", - "fanout": 6025, - "net": "g_tex_core[0].tex_core/rsp_buf/g_eb2.stream_buffer/mem_rsp_ready_15" + "fanout": 9440, + "net": "g_tex_core[0].tex_core/rsp_buf/g_eb2.stream_buffer/mem_rsp_ready_9" }, { "driver": "LUT2", @@ -136,42 +136,42 @@ }, { "driver": "BUFGCE", - "fanout": 3609, + "fanout": 3601, "net": "g_tex_core[0].tex_core/rsp_buf/g_eb2.stream_buffer/g_buffer.valid_in_r_reg_1" }, { "driver": "BUFGCE", - "fanout": 3568, - "net": "g_tex_core[0].tex_core/rsp_buf/g_eb2.stream_buffer/mem_bus_if\\\\.rsp_ready" + "fanout": 2960, + "net": "g_rtu_core[0].rtu_core/scheduler/ray_store/s1_valid_reg" }, { - "driver": "FDRE", - "fanout": 2901, - "net": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/g_req_queue_ro.req_queue/g_ebN.fifo_queue/g_depth_n.wr_ptr_r[0]" + "driver": "PORT", + "fanout": 2960, + "net": "reset" }, { - "driver": "FDRE", - "fanout": 2897, - "net": "g_tex_core[0].tex_core/tex_mem/mem_scheduler/g_req_queue_ro.req_queue/g_ebN.fifo_queue/g_depth_n.wr_ptr_r[1]" + "driver": "BUFGCE", + "fanout": 2297, + "net": "g_tex_core[0].tex_core/rsp_buf/g_eb2.stream_buffer/lsu_mem_if\\\\.rsp_ready" }, { - "driver": "BUFGCE", - "fanout": 2756, - "net": "g_rtu_core[0].rtu_core/scheduler/word_q" + "driver": "LUT5", + "fanout": 2274, + "net": "l2cache/g_cache.cache/g_banks[0].bank/reg_cmt/enable0" }, { - "driver": "PORT", - "fanout": 2706, - "net": "reset" + "driver": "LUT5", + "fanout": 2274, + "net": "l2cache/g_cache.cache/g_banks[1].bank/reg_cmt/enable0" }, { "driver": "LUT5", - "fanout": 2392, - "net": "l2cache/g_cache.cache/g_banks[0].bank/reg_cmt/enable0" + "fanout": 2274, + "net": "l2cache/g_cache.cache/g_banks[2].bank/reg_cmt/enable0" } ], - "lut": 202801, - "lutram": 10287, + "lut": 204871, + "lutram": 16721, "uram": 0, "wns_ns": 0.002 }