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/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/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 } 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..31b1951e13 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": { @@ -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": { 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/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/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/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) diff --git a/ci/testcases/asic_gate.yaml b/ci/testcases/asic_gate.yaml index 49a1f5d63b..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 @@ -33,14 +34,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 +124,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 +191,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 ---------------------------------------------------------------------- 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] 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 | |---|---|---| 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 // ============================================================================ 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)