From 9deacbee4c0e12f19065223dc810cbff70700544 Mon Sep 17 00:00:00 2001 From: tinebp Date: Wed, 30 Sep 2026 10:04:55 -0700 Subject: [PATCH 1/8] ci: fit the perf and parity gates in a hosted job; align SimX dispatch and DXA drain with the RTL The perf_gate and model_parity jobs timed out (run 36547574183). Since #416 the ALU/LSU/FPU block counts follow ISSUE_WIDTH, so every ISSUE_WIDTH=4 gate case elaborates a four-block model several times the size, and a hosted runner cannot build it within the job. Those cases now pin one block per unit; their cycles return to the pre-#416 values (wgmma-dxa-mcast nt4/nt16 reproduce them exactly). Two build costs amplified it: - rtlsim passed -O2 through CXXFLAGS, which lands after Verilator's per-class knobs and also optimizes the slow-path files. The level now goes through OPT_FAST/OPT_GLOBAL: wgmma-dxa 598 -> 224 CPU-min, same cycles. - blackbox built the driver with the caller's CONFIGS, then the run rebuilt it with the app's resolved CONFIGS. A runtime- target in each suite's common.mk builds it once, with the CONFIGS the run uses. SimX timing, both mismatches found on the DXA parity cases: - The dispatcher stepped through the issue-slot batches in fixed rotation, spending a cycle on every empty slot. VX_lane_dispatch grants the lowest batch holding an instruction and rests on the last when none does. - The DXA scatter drain (K-major, Flat, BlockMajor) gathered a word of elements per beat. VX_dxa_smem_wr writes one element per beat; DXA lmem_writes on wgmma-dxa now equal the RTL's. SimX against recorded RTL cycles, all 31 model_parity cases (RTL unchanged): wgmma-dxa 24.9% -> 5.5%, wgmma-sp-dxa 17.8% -> 5.2%, wgmma-dxa-mcast 9.8% -> 0.9%, sgemm2-fp16 4.5% -> 0.1%, wgmma-fedp2k-rs 4.0% -> 0.7%; every other case unchanged. The multicast case drops its known_issue and 10% tolerance, the sparse case its 10% tolerance; the two remaining known_issue reasons now state the measured gap. The wgmma-dxa residual is C-tile store drain at CTA turnover. Not run here: the full rtlsim model_parity and perf_gate gates. perf baselines are stale (fingerprint) and need --update-baselines. Co-Authored-By: Claude Opus 4.8 (1M context) --- ci/blackbox.sh | 9 +++-- ci/testcases/dxa.yaml | 49 +++++++++++++------------- ci/testcases/tensor.yaml | 12 +++++-- ci/testcases/tensor_wg.yaml | 20 ++++++++--- docs/designs/continuous_integration.md | 11 +++--- sim/rtlsim/Makefile | 7 +++- sim/simx/dispatcher.cpp | 15 ++++++-- sim/simx/dxa/dxa_core.cpp | 40 ++++++++++----------- tests/graphics/common.mk | 5 +++ tests/hip/common.mk | 5 +++ tests/mpi/common.mk | 5 +++ tests/opencl/common.mk | 5 +++ tests/regression/common.mk | 5 +++ tests/vulkan/common.mk | 5 +++ 14 files changed, 129 insertions(+), 64 deletions(-) diff --git a/ci/blackbox.sh b/ci/blackbox.sh index 963c1a80b9..f4072e9ada 100755 --- a/ci/blackbox.sh +++ b/ci/blackbox.sh @@ -130,14 +130,17 @@ build_driver() { [ $VCD -eq 1 ] && cmd_opts=$(add_option "$cmd_opts" "VCD=1") [ $SAIF -eq 1 ] && cmd_opts=$(add_option "$cmd_opts" "SAIF=1") [ -n "$TARGET" ] && cmd_opts=$(add_option "$cmd_opts" "TARGET=$TARGET") - [ $TEMPBUILD -eq 1 ] && cmd_opts=$(add_option "$cmd_opts" "DESTDIR=\"$TEMPDIR\"") + [ $TEMPBUILD -eq 1 ] && cmd_opts=$(add_option "$cmd_opts" "VORTEX_RT_LIB=\"$TEMPDIR\"") [ -n "$CONFIGS" ] && cmd_opts=$(add_option "$cmd_opts" "CONFIGS=\"$CONFIGS\"") - cmd_opts=$(add_option "$cmd_opts" "make -C $DRIVER_PATH > /dev/null") + # Through the app's Makefile, not the driver's: an app resolves its own + # CONFIGS defaults, and a driver built from the caller's CONFIGS alone is a + # different model that the run then discards and rebuilds. + cmd_opts=$(add_option "$cmd_opts" "make -C \"$APP_PATH\" runtime-$DRIVER > /dev/null") echo "Running: $cmd_opts" eval "$cmd_opts" status=$? if [ $status -ne 0 ]; then - echo "Error building driver: $DRIVER_PATH" + echo "Error building driver: $DRIVER (via $APP_PATH)" exit $status fi } diff --git a/ci/testcases/dxa.yaml b/ci/testcases/dxa.yaml index 2751838e92..da6b39bc7f 100644 --- a/ci/testcases/dxa.yaml +++ b/ci/testcases/dxa.yaml @@ -507,15 +507,17 @@ tests: threads: [4, 16] - id: model_parity-wgmma-dxa check: model_parity - # Instructions match exactly; cycles diverge ~6% (RTL faster). The gap is - # insensitive to every issue/commit/execution-window mechanism the SimX - # timing model carries, which points at the DXA producer path's own - # timing. Tracked with its wgmma siblings as DXA-path follow-up. - known_issue: "SimX over-serializes the DXA-fed warp-group pipeline (instrs match; cycles ~6%)" + # Instructions and DXA transfer timing match; cycles diverge ~5% (SimX + # slower), all of it at CTA turnover: SimX drains the C-tile stores later + # than the RTL (store commit point, dispatch queue placement and bank-0 + # store-miss admission). + known_issue: "SimX drains the C-tile stores at CTA turnover slower than the RTL (instrs match; cycles ~5%)" via: blackbox app: sgemm_tcu_wg_dxa args: -m 128 -n 128 -k 64 - configs: -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DWGMMA_NRC=16 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 # Not warp-width swept: a warp-group MMA is defined over a 32-lane warp. shape: {threads: 32, warps: 8} - id: perf_gate-wgmma-dxa @@ -523,51 +525,50 @@ tests: via: blackbox app: sgemm_tcu_wg_dxa args: -m 128 -n 128 -k 64 - configs: -DVX_CFG_ISSUE_WIDTH=4 -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_TCU_WGMMA_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DWGMMA_NRC=16 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 # Not warp-width swept: a warp-group MMA is defined over a 32-lane warp. shape: {threads: 32, warps: 8} # Sparse WGMMA consumer fed by a DXA producer pipeline. - id: model_parity-wgmma-sp-dxa check: model_parity - # Instructions match exactly; cycles diverge ~15% (RTL slower). The - # contended issue-path serializers that closed core:parity-sgemm-mc shave - # only ~2 points here, and the residual is insensitive to every - # issue/commit/execution-window mechanism, which points at the DXA - # producer path's own timing under the 32-thread ISSUE_WIDTH=4 shape. - # Tracked with its wgmma siblings as DXA-path follow-up. - known_issue: "SimX under-models the DXA producer path for 32-thread warp groups (instrs match; cycles ~15%)" - tolerance: 0.10 + # Instructions match; cycles diverge ~5% (SimX faster). Not yet localized. + known_issue: "SimX runs the sparse DXA-fed warp-group kernel faster than the RTL (instrs match; cycles ~5%)" via: blackbox app: sgemm_tcu_wg_sp_dxa args: -m 128 -n 128 -k 128 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 shape: {threads: 32, warps: 8} - id: perf_gate-wgmma-sp-dxa check: perf_gate via: blackbox app: sgemm_tcu_wg_sp_dxa args: -m 128 -n 128 -k 128 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 shape: {threads: 32, warps: 8} # WGMMA + multicast DXA: the large shared operand is broadcast once. K=64 for # the same reference-tolerance reason as perf_gate-wgmma-dxa. - id: model_parity-wgmma-dxa-mcast check: model_parity - # Same family as model_parity-wgmma-sp-dxa: instrs match, cycles diverge - # (~12%, RTL slower), insensitive to the issue/commit/execution-window - # mechanisms. Tracked as DXA-path follow-up. - known_issue: "SimX under-models the multicast DXA producer path (instrs match; cycles ~12%)" - tolerance: 0.10 via: blackbox app: sgemm_tcu_wg_dxa_mcast args: -m 128 -n 128 -k 64 - configs: -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 - id: perf_gate-wgmma-dxa-mcast check: perf_gate via: blackbox app: sgemm_tcu_wg_dxa_mcast args: -m 128 -n 128 -k 64 - configs: -DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 shape: threads: [4, 16] \ No newline at end of file diff --git a/ci/testcases/tensor.yaml b/ci/testcases/tensor.yaml index 34bd128da8..0b767486f8 100644 --- a/ci/testcases/tensor.yaml +++ b/ci/testcases/tensor.yaml @@ -187,7 +187,9 @@ tests: via: blackbox app: sgemm_tcu args: -m 128 -n 128 -k 128 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 # TCU fp16 multi-core: NT=NW=8 across 2 cores (shared L2). # Not warp-width swept: the tensor-core fragment layout is defined over a fixed @@ -216,13 +218,17 @@ tests: via: blackbox app: sgemm2_tcu args: -m 128 -n 128 -k 64 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 - id: perf_gate-sgemm2-fp16 check: perf_gate via: blackbox app: sgemm2_tcu args: -m 128 -n 128 -k 64 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 # Not warp-width swept: the tensor-core fragment layout is defined over a fixed # lane count, which the configs pin. \ No newline at end of file diff --git a/ci/testcases/tensor_wg.yaml b/ci/testcases/tensor_wg.yaml index 52514adfff..fd3449444a 100644 --- a/ci/testcases/tensor_wg.yaml +++ b/ci/testcases/tensor_wg.yaml @@ -196,7 +196,9 @@ tests: check: model_parity via: blackbox app: sgemm_tcu_wg - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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_INT8_ENABLE -DITYPE=int8 -DOTYPE=int32 -DWGMMA_NRC=32 # Dense WGMMA fp16 SS datapath, cycles-vs-baseline. @@ -205,14 +207,18 @@ tests: via: blackbox app: sgemm_tcu_wg args: -m 128 -n 128 -k 128 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 - id: perf_gate-wgmma-fedp2k-rs check: perf_gate via: blackbox app: sgemm_tcu_wg args: -m64 -n64 -k64 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 - id: model_parity-wgmma-fedp2k-rs @@ -222,7 +228,9 @@ tests: via: blackbox app: sgemm_tcu_wg args: -m64 -n64 -k64 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 -DVX_CFG_TCU_TYPE_DPI -DITYPE=fp16 -DOTYPE=fp32 -DWGMMA_NRC=8 -DWGMMA_RS # Cooperative sparse WGMMA (default fp16 operands, 2:4 structured sparsity). @@ -233,7 +241,9 @@ tests: via: blackbox app: sgemm_tcu_wg_sp args: -m 128 -n 128 -k 128 - 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 + # One execution block per unit: the four-block model ISSUE_WIDTH=4 implies is + # several times the size, and a hosted CI runner cannot build it in a job. + 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 # Not warp-width swept: the tensor-core fragment layout is defined over a fixed # lane count, which the configs pin. \ No newline at end of file diff --git a/docs/designs/continuous_integration.md b/docs/designs/continuous_integration.md index 23ca5d6e35..7bc048fd76 100644 --- a/docs/designs/continuous_integration.md +++ b/docs/designs/continuous_integration.md @@ -202,12 +202,11 @@ A check is an assertion across runs rather than a clean exit. | tier | `full` | `full` | | pinned to | the rtlsim driver, which elaborates the RTL | the rtlsim driver | -**A check is a marker, never a file or a category.** Each check gets its own -cells — one per suite it has cases in — and every category cell excludes the -check markers, so a check case runs exactly once. One cell per suite rather -than one for the whole catalog, because every case is a simulator build of -its own and a single cell outgrew a hosted runner's job limit. The workflow -reads the check list from `testcase.py checks` rather than holding a copy. +**A check is a marker, never a file or a category.** Each check gets one cell +of its own — `-m " and rtlsim"` sweeps every such case catalog-wide — +and every category cell excludes the check markers, so a check case runs +exactly once. The workflow reads the check list from `testcase.py checks` +rather than holding a copy. #### `model_parity` diff --git a/sim/rtlsim/Makefile b/sim/rtlsim/Makefile index b935e8458a..9856159b26 100644 --- a/sim/rtlsim/Makefile +++ b/sim/rtlsim/Makefile @@ -169,7 +169,12 @@ ifdef DEBUG CXXFLAGS += -g -O0 $(DBG_FLAGS) else VL_FLAGS += -DNDEBUG - CXXFLAGS += -O2 -DNDEBUG + CXXFLAGS += -DNDEBUG + # The level goes through Verilator's per-class knobs, not CXXFLAGS: a level + # in CXXFLAGS lands after them and so also optimizes the slow-path files + # (constructors, reset, initialization), which run once and are most of the + # compile time of a large model. + VL_FLAGS += -MAKEFLAGS OPT_FAST=-O2 -MAKEFLAGS OPT_GLOBAL=-O2 endif # Enable perf counters diff --git a/sim/simx/dispatcher.cpp b/sim/simx/dispatcher.cpp index 30f6109f91..273c6747cc 100644 --- a/sim/simx/dispatcher.cpp +++ b/sim/simx/dispatcher.cpp @@ -41,6 +41,16 @@ void Dispatcher::on_reset() { } void Dispatcher::on_tick() { + // Batches holding an instruction this cycle, sampled before any is popped. + uint32_t valid_batches = 0; + if (num_blocks_ != 1) { + for (uint32_t i = 0; i < VX_CFG_ISSUE_WIDTH; ++i) { + if (!Inputs.at(i).empty()) { + valid_batches |= 1u << (i / block_size_); + } + } + } + // process inputs uint32_t block_sent = 0; for (uint32_t b = 0; b < block_size_; ++b) { @@ -126,8 +136,9 @@ void Dispatcher::on_tick() { // advance to next batch once all blocks in the current batch have been processed if (block_sent == block_size_) { - // round-robin batch selection - batch_idx_ = (batch_idx_ + 1) % num_blocks_; + // Priority grant to the lowest batch with an instruction: an empty issue + // slot costs no dispatch cycle. With none, the grant rests on the last. + batch_idx_ = valid_batches ? __builtin_ctz(valid_batches) : (num_blocks_ - 1); for (auto& bp : block_pids_) { bp = 0; } diff --git a/sim/simx/dxa/dxa_core.cpp b/sim/simx/dxa/dxa_core.cpp index 3f51c3254e..eb9b42dbb3 100644 --- a/sim/simx/dxa/dxa_core.cpp +++ b/sim/simx/dxa/dxa_core.cpp @@ -67,11 +67,9 @@ class DxaCore::Impl { // ── Per-line work item produced by addr_gen, consumed by gmem_req ─── // Each entry = ONE GMEM cache-line read. Row-major writes contribute it to a - // single contiguous LMEM word write; a K-major transposing load fans it out - // to `km_num_elems` strided per-element destinations, which smem_wr coalesces - // back into full byte-masked block writes (one read → bank-parallel scatter, - // never re-reading a line per element — TMA-style, matching the hardware - // address generator, which reads per cache line). + // single contiguous LMEM word write; a K-major or tiled load fans it out to + // `km_num_elems` per-element destinations, which smem_wr writes one element + // per beat (one read per cache line, never re-read per element). struct LineWork { uint64_t gmem_cl_addr; // CL-aligned global address uint64_t smem_word_addr; // word-aligned SMEM byte address (element 0) @@ -81,6 +79,7 @@ class DxaCore::Impl { uint32_t cfill; // OOB fill value (lane-replicated) bool oob; // skip GMEM read; use cfill bool last; // last work item of the transfer + bool scatter; // K-major/tiled: one element written per beat uint32_t km_num_elems; // K-major scatter fan-out (1 = contiguous write) uint32_t km_lane_stride; // SMEM byte stride between scattered elements uint8_t dest_layout; // DestLayout (Flat/BlockMajor use tiled_dest_elem) @@ -509,12 +508,12 @@ class DxaCore::Impl { lw.cl_byte_offset = cl_off; lw.smem_byte_offset = s_off; // Row-major: one contiguous write of the whole span. Scatter (K-major - // or tiled Flat/BlockMajor): each element is its own write (elem_bytes), - // gathered into block writes by smem_wr. + // or tiled Flat/BlockMajor): each element is its own write (elem_bytes). lw.valid_length = scatter ? elem_bytes : gspan; lw.cfill = cfill; lw.oob = elem_oob; lw.last = false; + lw.scatter = scatter; lw.km_num_elems = scatter ? num_elems : 1; lw.km_lane_stride = dest_kmajor ? per_lane_stride_bytes : 0; lw.dest_layout = uint8_t(layout); @@ -659,14 +658,11 @@ class DxaCore::Impl { const LineWork& lw = s.work; - // One cache-line read fans out to its scattered writes. The K-major - // elements landing in the same SMEM block are written TOGETHER in one - // byte-masked block write: the per-core LMEM port accepts a full - // VX_CFG_MEM_BLOCK_SIZE word per cycle (banked), exactly like the warp - // array's transpose — so the engine drains at SMEM bandwidth, not one - // element per cycle (TMA-style full-bandwidth scatter). Cursors: - // km_elem_idx advances one block per beat; mc_cta_idx replays the whole - // group to each multicast receiver. + // One cache-line read fans out to its scattered writes, one element per + // beat: the scatter destinations are not contiguous, so each element is + // its own byte-masked write. Cursors: km_elem_idx advances one element + // per beat; mc_cta_idx replays the whole group to each multicast + // receiver. const uint32_t num_elems = lw.km_num_elems ? lw.km_num_elems : 1; const uint32_t e0 = w.km_elem_idx; const uint32_t wlen = lw.valid_length; @@ -730,10 +726,11 @@ class DxaCore::Impl { auto blk = make_mem_block(); const uint32_t pat = lw.cfill; - // Gather every scatter element that falls in this block (and, for rows - // narrower than a block, this row) into one write. + // A contiguous span is one write; a scatter span writes element e0. + const bool beat_scatter = lw.scatter; + const uint32_t e_end = beat_scatter ? (e0 + 1) : num_elems; uint32_t ee = e0; - for (; ee < num_elems; ++ee) { + for (; ee < e_end; ++ee) { uint64_t dest_byte = dest_byte_of(ee); if ((dest_byte & ~uint64_t(kLmemWordSize - 1)) != dword) break; if ((dest_byte & kRowMask) != row) { @@ -754,7 +751,7 @@ class DxaCore::Impl { req.byteen = byteen; req.data = blk; - // notify_done on the LAST block write of the transfer — when the gather + // notify_done on the LAST block write of the transfer — when the write // reached the last scatter element of the last work item, per receiver. bool is_last_elem = (ee == num_elems); bool is_last_work = lw.last; @@ -773,7 +770,7 @@ class DxaCore::Impl { ++perf_stats_.lmem_writes; } - // Advance scatter cursor (to first ungathered element); then multicast + // Advance scatter cursor (to the next unwritten element); then multicast // cursor; then release the slot. if (!is_last_elem) { w.km_elem_idx = ee; @@ -792,6 +789,9 @@ class DxaCore::Impl { w.drain_slot = UINT32_MAX; } } + if (beat_scatter) { + break; // a scatter beat writes one element + } } // per-tick row-beat emission loop } diff --git a/tests/graphics/common.mk b/tests/graphics/common.mk index 2e899cac7c..8838140145 100644 --- a/tests/graphics/common.mk +++ b/tests/graphics/common.mk @@ -193,6 +193,11 @@ $(RT_LIB_DIR)/libvortex.so: $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/stub HOST_ARCH=$(HOST_ARCH) DESTDIR=$(VORTEX_RT_LIB) endif +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + run-simx: $(PROJECT) kernel.vxbin $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/simx DESTDIR=$(VORTEX_RT_LIB) LD_LIBRARY_PATH=$(VORTEX_RT_LIB):$(LD_LIBRARY_PATH) VORTEX_DRIVER=simx ./$(PROJECT) $(OPTS) diff --git a/tests/hip/common.mk b/tests/hip/common.mk index 2266a8a94d..0b315896d2 100644 --- a/tests/hip/common.mk +++ b/tests/hip/common.mk @@ -142,6 +142,11 @@ $(VORTEX_RT_LIB)/libvortex.so: FORCE $(PROJECT): $(SRCS) common.h $(VORTEX_KN_PATH)/libvortex2.a $(VORTEX_RT_LIB)/libvortex.so HIP_CLANG_PATH=$(HIP_CLANG_PATH) LD_LIBRARY_PATH=$(LLVM_PATH)/lib:$(LD_LIBRARY_PATH) $(HIPCC) $(HIPCC_FLAGS) -I. $< -o $@ +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + run-simx: $(PROJECT) $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/simx DESTDIR=$(VORTEX_RT_LIB) $(HIP_OCL_ENV) LD_LIBRARY_PATH=$(OCL_ICD_LIB_DIR):$(CHIPSTAR_PATH)/lib:$(POCL_PATH)/lib:$(VORTEX_RT_LIB):$(LLVM_PATH)/lib:$(LD_LIBRARY_PATH) $(POCL_CC_FLAGS) VORTEX_DRIVER=simx ./$(PROJECT) $(OPTS) diff --git a/tests/mpi/common.mk b/tests/mpi/common.mk index 101ad4656b..1a0c244260 100644 --- a/tests/mpi/common.mk +++ b/tests/mpi/common.mk @@ -193,6 +193,11 @@ $(RT_LIB_DIR)/libvortex.so: $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/stub HOST_ARCH=$(HOST_ARCH) DESTDIR=$(VORTEX_RT_LIB) endif +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + run-simx: $(PROJECT) kernel.vxbin $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/simx DESTDIR=$(VORTEX_RT_LIB) LD_LIBRARY_PATH=$(VORTEX_RT_LIB):$(LD_LIBRARY_PATH) VORTEX_DRIVER=simx ./$(PROJECT) $(OPTS) diff --git a/tests/opencl/common.mk b/tests/opencl/common.mk index 8b4e2c8e19..e35673e00e 100644 --- a/tests/opencl/common.mk +++ b/tests/opencl/common.mk @@ -122,6 +122,11 @@ $(PROJECT).host: $(OBJS) run-gpu: $(PROJECT).host $(KERNEL_SRCS) ./$(PROJECT).host $(OPTS) +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + run-simx: $(PROJECT) $(KERNEL_SRCS) $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/simx DESTDIR=$(VORTEX_RT_LIB) LD_LIBRARY_PATH=$(OCL_ICD_LIB_DIR):$(POCL_PATH)/lib:$(VORTEX_RT_LIB):$(LLVM_PATH)/lib:$(LD_LIBRARY_PATH) $(POCL_CC_FLAGS) OCL_ICD_VENDORS=$(OCL_ICD_VENDORS) VORTEX_DRIVER=simx ./$(PROJECT) $(OPTS) diff --git a/tests/regression/common.mk b/tests/regression/common.mk index 47e4cfeff9..ed7a588772 100644 --- a/tests/regression/common.mk +++ b/tests/regression/common.mk @@ -223,6 +223,11 @@ $(RT_LIB_DIR)/libvortex.so: $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/stub HOST_ARCH=$(HOST_ARCH) DESTDIR=$(VORTEX_RT_LIB) endif +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + run-simx: $(PROJECT) kernel.vxbin $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/simx DESTDIR=$(VORTEX_RT_LIB) LD_LIBRARY_PATH=$(VORTEX_RT_LIB):$(LD_LIBRARY_PATH) VORTEX_DRIVER=simx ./$(PROJECT) $(OPTS) diff --git a/tests/vulkan/common.mk b/tests/vulkan/common.mk index 8b72b2746c..a3a6453258 100644 --- a/tests/vulkan/common.mk +++ b/tests/vulkan/common.mk @@ -149,6 +149,11 @@ $(PROJECT): $(SRCS) $(VORTEX_RT_LIB)/libvortex.so: FORCE $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/stub DESTDIR=$(VORTEX_RT_LIB) +# The runtime for one driver, built with the CONFIGS this app resolved, so a +# caller that builds ahead of the run builds the model the run will use. +runtime-%: + $(RUNTIME_ARGS) $(MAKE) -C $(VORTEX_RT_SRC)/$* DESTDIR=$(VORTEX_RT_LIB) + # `run` defaults to the SimX backend; explicit recipes select simx / rtlsim # / opae / xrt. vortexpipe is backend-agnostic — same .vxbin, the stub # libvortex.so dlopens libvortex-.so at runtime. From 5e65cbaf010f3eb8b22457bfcc8541488a473a46 Mon Sep 17 00:00:00 2001 From: tinebp Date: Wed, 30 Sep 2026 17:57:27 -0700 Subject: [PATCH 2/8] simx: hold FPU tags until commit and gate issue on going-full FUs as the RTL does SGEMM on two cores (model_parity sgemm-mc) ran 10.5% slower in SimX than in the RTL. The kernel's loop is a dependent eight-FMA chain with thirteen loads, so it is bound by the FPU's two tags. Three SimX timing paths diverged from the RTL: - Issue: VX_scoreboard drops a warp from arbitration while its target FU's dispatch queue is going-full. SimX only de-prioritized such warps and let them all through when every ready warp targeted a congested FU. Excess FPU ops (7 issued and not yet in the FPU, against the RTL's 4) filled the two shared operand collectors and blocked loads and ALU ops behind them. The warps fell into lockstep and their load writebacks committed in bursts of 17-40 beats that starved the lower-priority FPU. The gate is now hard. The suppress-mask arbiter overloads lose their last caller and are removed. - FPU: VX_fpu_unit holds each tag until the result fires into the commit arbiter (no buffer between them at full bandwidth). SimX freed tags on a timer, buffered extra results and added two cycles to every latency. The response port now has one slot per tag, so taking a result off it frees the tag. At full bandwidth the commit arbiter reads the port directly. With partial bandwidth the lane-gather stage does. The latencies are the configured unit latencies. - Crossbar: TxCrossBar ignored its arbiter type and always granted the lowest input. Round-robin outputs now rotate from the last accepted input, reset to the last input, as VX_stream_xbar / VX_rr_arbiter do. The cache crossbars (round-robin) change; local memory (priority) does not. SimX against the recorded RTL cycles, all 31 model_parity cases: seven were over 5%, three are now. sgemm-mc 10.49% -> 2.06%, sgemmx 3.93 -> 0.53, rt_raycast 6.27 -> 2.83, fp16 6.28 -> 3.32, fp16-mc 4.07 -> 1.47, wgmma-sp-dxa 5.15 -> 4.21. Three got worse, all from the issue gate: sgemv 7.23 -> 9.59, wgmma-dxa 5.46 -> 8.53, mxfp8 0.82 -> 3.36. The gate matches the RTL, so these expose other SimX timing errors it was offsetting. They are not localized yet. Catalog: sgemmx and fp16 return to the 5% default and wgmma-sp-dxa loses its known_issue. sgemv and copy-mcast keep their existing 10% tolerance, and the wgmma-dxa reason states its measured gap. Verified: SimX-only sweep of all 31 model_parity cases against recorded RTL cycles (the RTL is unchanged). The final channel-based FPU wiring reproduces the swept cycles exactly on sgemm-mc, sgemv, fp16 and wgmma-dxa. sgemmx moves 2 cycles because the event wheel's push_back does not preserve FIFO order (fixed separately). FP-heavy apps pass in both FPU gather modes. Not run: the full pytest model_parity suite with its rtlsim legs. perf_gate runs on rtlsim only and cannot move. Co-Authored-By: Claude Opus 4.8 (1M context) --- ci/testcases/core.yaml | 11 ++------ ci/testcases/dxa.yaml | 16 +++++------ ci/testcases/tensor.yaml | 5 ---- docs/designs/floating_point_unit.md | 28 ++++++++++++------ sim/simx/core.cpp | 38 ++++++++++++++----------- sim/simx/fpu_unit.cpp | 39 ++++++++++--------------- sim/simx/fpu_unit.h | 14 ++++----- sim/simx/types.h | 44 +++++++++++++---------------- 8 files changed, 90 insertions(+), 105 deletions(-) diff --git a/ci/testcases/core.yaml b/ci/testcases/core.yaml index 92532beef2..4a2bf69000 100644 --- a/ci/testcases/core.yaml +++ b/ci/testcases/core.yaml @@ -146,10 +146,9 @@ tests: # SGEMV: FPU + streaming matrix-vector. - id: parity-sgemv check: model_parity - # Instrs match exactly; cycles diverge ~7% (RTL slower) in the high-miss - # streaming regime, where request-arrival timing at the cache shifts the - # hit/miss split between the models even though isolated round-trips match. - tolerance: 0.10 # residual streaming-regime divergence (instrs match; cycles ~7%) + # Instrs match exactly; cycles diverge ~9.6% (SimX slower). Not yet + # localized. + tolerance: 0.10 # instrs match; cycles ~9.6% via: blackbox app: sgemv args: -m512 -n512 @@ -168,10 +167,6 @@ tests: # SGEMM with local-memory tiling (kernel-managed shared memory). - id: parity-sgemmx check: model_parity - # Instrs match exactly; cycles diverge ~10% (RTL slower): the local-memory - # store/writeback serialization keeps a residual the execution-window - # models do not fully close. - tolerance: 0.10 via: blackbox app: sgemmx args: -n128 diff --git a/ci/testcases/dxa.yaml b/ci/testcases/dxa.yaml index da6b39bc7f..2941955993 100644 --- a/ci/testcases/dxa.yaml +++ b/ci/testcases/dxa.yaml @@ -459,8 +459,8 @@ tests: threads: [4, 16] - id: model_parity-copy-mcast check: model_parity - # Same cause as model_parity-copy: instrs match, cycles diverge (~9%). - tolerance: 0.10 # SimX does not model the decoupled LSU pending pool (instrs match; cycles ~9%) + # Instrs match; cycles diverge ~6.8% (SimX slower). Not yet localized. + tolerance: 0.10 # instrs match; cycles ~6.8% via: blackbox app: dxa_copy_mcast configs: -DVX_CFG_EXT_DXA_ENABLE @@ -507,11 +507,11 @@ tests: threads: [4, 16] - id: model_parity-wgmma-dxa check: model_parity - # Instructions and DXA transfer timing match; cycles diverge ~5% (SimX - # slower), all of it at CTA turnover: SimX drains the C-tile stores later - # than the RTL (store commit point, dispatch queue placement and bank-0 - # store-miss admission). - known_issue: "SimX drains the C-tile stores at CTA turnover slower than the RTL (instrs match; cycles ~5%)" + # Instructions and DXA transfer timing match; cycles diverge ~8.5% (SimX + # slower), at CTA turnover: SimX drains the C-tile stores later than the RTL + # (store commit point, dispatch queue placement and bank-0 store-miss + # admission). + known_issue: "SimX drains the C-tile stores at CTA turnover slower than the RTL (instrs match; cycles ~8.5%)" via: blackbox app: sgemm_tcu_wg_dxa args: -m 128 -n 128 -k 64 @@ -533,8 +533,6 @@ tests: # Sparse WGMMA consumer fed by a DXA producer pipeline. - id: model_parity-wgmma-sp-dxa check: model_parity - # Instructions match; cycles diverge ~5% (SimX faster). Not yet localized. - known_issue: "SimX runs the sparse DXA-fed warp-group kernel faster than the RTL (instrs match; cycles ~5%)" via: blackbox app: sgemm_tcu_wg_sp_dxa args: -m 128 -n 128 -k 128 diff --git a/ci/testcases/tensor.yaml b/ci/testcases/tensor.yaml index 0b767486f8..062a8dd331 100644 --- a/ci/testcases/tensor.yaml +++ b/ci/testcases/tensor.yaml @@ -163,11 +163,6 @@ tests: # config on both drivers, cycles must agree within tolerance (see core.yaml). - id: model_parity-fp16 check: model_parity - # Execution passes and retired instructions match exactly; only cycles - # diverge (~6%, just over the 5% gate). SimX does not model the decoupled LSU - # pending pool that feeds the TCU operand loads. Parity divergence only; - # pending SimX LSU timing alignment. - tolerance: 0.10 # SimX does not model the decoupled LSU pending pool (instrs match; cycles ~6%) via: blackbox app: sgemm_tcu configs: -DVX_CFG_NUM_THREADS=8 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32 diff --git a/docs/designs/floating_point_unit.md b/docs/designs/floating_point_unit.md index 0ced9c4548..65cd877c8a 100644 --- a/docs/designs/floating_point_unit.md +++ b/docs/designs/floating_point_unit.md @@ -402,18 +402,28 @@ The Vivado and Quartus rows are the vendor operators' fixed latencies. ## 11. SimX model [`FpuUnit`](../../sim/simx/fpu_unit.cpp) computes results with the -`rvfloats` soft-float library and models timing as the operation's latency -plus a fixed 2 cycles for the front end: +`rvfloats` soft-float library and models timing with the configured unit +latencies: | operation class | modeled latency | |---|---| -| FMA | `VX_CFG_FMA_LATENCY` + 2 | -| DIV, SQRT | `VX_CFG_FDIV_LATENCY` + 2, `VX_CFG_FSQRT_LATENCY` + 2 | -| CVT | `VX_CFG_FCVT_LATENCY` + 2 | -| NCP | 2 + 2 | - -Its output capacity is the tag-store depth plus one result-skid entry, which -bounds the operations in flight the same way the RTL's tag store does. +| FMA | `VX_CFG_FMA_LATENCY` | +| DIV, SQRT | `VX_CFG_FDIV_LATENCY`, `VX_CFG_FSQRT_LATENCY` | +| CVT | `VX_CFG_FCVT_LATENCY` | +| NCP | `VX_CFG_FNCP_LATENCY` | + +Its response port has one slot per tag (`VX_CFG_FPU_QUEUE_SIZE`). As in the +RTL, an operation holds its tag from acceptance until its result is taken off +that port, so the tags, not the pipeline depth, bound the operations in +flight, and a result waiting on the commit arbiter keeps its tag: + +- At full bandwidth (`NUM_FPU_BLOCKS == ISSUE_WIDTH` and + `NUM_FPU_LANES == SIMD_WIDTH`) the commit arbiter reads the response port + directly and frees the tag when it grants the result. The result also takes + the one registered stage the other units cross into commit, so an FMA + commits `VX_CFG_FMA_LATENCY + 1` cycles after it enters. +- At partial bandwidth the lane-gather stage takes the result off the port and + frees the tag. --- diff --git a/sim/simx/core.cpp b/sim/simx/core.cpp index 61e80e974c..1e447cec6d 100644 --- a/sim/simx/core.cpp +++ b/sim/simx/core.cpp @@ -539,7 +539,6 @@ class Core::Impl { for (uint32_t iw = 0; iw < VX_CFG_ISSUE_WIDTH; ++iw) { bool any_scrb_blocked = false; BitVector<> ready_set(PER_ISSUE_WARPS); - BitVector<> suppress_set(PER_ISSUE_WARPS); for (uint32_t w = 0; w < PER_ISSUE_WARPS; ++w) { uint32_t wid = w * VX_CFG_ISSUE_WIDTH + iw; auto& ibuffer = ibuffers_.at(wid); @@ -575,6 +574,12 @@ class Core::Impl { if (fu_locked_.at(iw).test(fu) && uop_fu_lock) { continue; // blocked by FU lock } + // FU dispatch queue going-full: the warp does not request. Credits also + // count ops still in operand collection; the one-slot guard band keeps + // an issued op from blocking the shared operand path. + if (fu_credits_.at(iw).at(fu) >= VX_CFG_DISPATCH_QUEUE_SIZE - 1) { + continue; + } #ifdef VX_CFG_EXT_RTU_ENABLE // A TRACE macro must hold a ray-pool slot before its head uop enters // the SFU, or it stalls at the head of that unit's queue and starves @@ -586,26 +591,12 @@ class Core::Impl { } #endif ready_set.set(w); // mark instruction as ready - // suppress warps whose target FU dispatch queue is going-full. Credit - // based: spent at issue, returned at FU accept, so it counts in-flight - // ops still in operand collection (like the hardware scoreboard), not just - // what has already reached the queue. - if (fu_credits_.at(iw).at(fu) >= VX_CFG_DISPATCH_QUEUE_SIZE - 1) { - suppress_set.set(w); - } } } if (ready_set.any()) { - // Only suppress when at least one warp can issue to a free FU; - // otherwise let all warps through so the pipeline absorbs transient stalls. - BitVector<> eff_suppress(PER_ISSUE_WARPS); - auto unsuppressed = ready_set & ~suppress_set; - if (unsuppressed.any()) { - eff_suppress = suppress_set; - } // select one instruction from ready set - auto w = ibuffer_arbs_.at(iw).grant(ready_set, eff_suppress); + auto w = ibuffer_arbs_.at(iw).grant(ready_set); uint32_t wid = w * VX_CFG_ISSUE_WIDTH + iw; auto& ibuffer = ibuffers_.at(wid); auto trace = ibuffer->peek(); @@ -723,6 +714,8 @@ class Core::Impl { // dispatcher aggregation, so we recover it from the warp id and try_send // into the matching commit queue. for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { + if (fu == (uint32_t)FUType::FPU && !FpuUnit::kGather) + continue; // commit reads the FPU response ports directly auto& func_unit = func_units_.at(fu); uint32_t nb = func_unit->num_blocks(); for (uint32_t b = 0; b < nb; ++b) { @@ -756,16 +749,27 @@ class Core::Impl { }; for (uint32_t iw = 0; iw < VX_CFG_ISSUE_WIDTH; ++iw) { SimChannel* granted = nullptr; + FUType granted_fu = FUType::ALU; for (auto fu : kCommitPrio) { - auto& queue = *commit_queues_.at(iw).at((uint32_t)fu); + // Without a lane-gather stage the FPU response port, one block per + // issue slot, feeds the arbiter: a result holds its FPU tag until + // granted. + auto& queue = (fu == FUType::FPU && !FpuUnit::kGather) + ? func_units_.at((uint32_t)fu)->output(iw) + : *commit_queues_.at(iw).at((uint32_t)fu); if (!queue.empty()) { granted = &queue; + granted_fu = fu; break; } } if (!granted) continue; auto trace = granted->peek(); + if (granted_fu == FUType::FPU && !FpuUnit::kGather + && trace->eop && trace->resume_warp) { + scheduler_->resume(trace->wid); + } // advance to commit stage DT(3, simobject_->name() << "-pipeline commit: " << *trace); diff --git a/sim/simx/fpu_unit.cpp b/sim/simx/fpu_unit.cpp index ef9531b01a..270f7b8a3f 100644 --- a/sim/simx/fpu_unit.cpp +++ b/sim/simx/fpu_unit.cpp @@ -40,14 +40,14 @@ inline int64_t check_boxing(int64_t a) { } FpuUnit::FpuUnit(const SimContext& ctx, const char* name, Core* core) - // Output capacity covers the tag-bounded operations in the datapath - // plus one result-skid entry ahead of commit. - : FuncUnit(ctx, name, core, VX_CFG_FPU_QUEUE_SIZE + 1) + // One output slot per tag: an operation holds its slot from acceptance + // until its result is taken off the response port, however deep the + // pipelines are. + : FuncUnit(ctx, name, core, VX_CFG_FPU_QUEUE_SIZE) {} uint32_t FpuUnit::latency_of(const instr_trace_t* trace) const { auto fpu_type = std::get(trace->op_type); - const uint32_t delay = 2; switch (fpu_type) { case FpuType::FCMP: case FpuType::FSGNJ: @@ -55,7 +55,7 @@ uint32_t FpuUnit::latency_of(const instr_trace_t* trace) const { case FpuType::FMVXW: case FpuType::FMVWX: case FpuType::FMINMAX: - return 2+delay; + return VX_CFG_FNCP_LATENCY; case FpuType::FADD: case FpuType::FSUB: case FpuType::FMUL: @@ -63,15 +63,15 @@ uint32_t FpuUnit::latency_of(const instr_trace_t* trace) const { case FpuType::FMSUB: case FpuType::FNMADD: case FpuType::FNMSUB: - return VX_CFG_FMA_LATENCY+delay; + return VX_CFG_FMA_LATENCY; case FpuType::FDIV: - return VX_CFG_FDIV_LATENCY+delay; + return VX_CFG_FDIV_LATENCY; case FpuType::FSQRT: - return VX_CFG_FSQRT_LATENCY+delay; + return VX_CFG_FSQRT_LATENCY; case FpuType::F2I: case FpuType::I2F: case FpuType::F2F: - return VX_CFG_FCVT_LATENCY+delay; + return VX_CFG_FCVT_LATENCY; default: std::abort(); } @@ -376,28 +376,19 @@ void FpuUnit::execute(instr_trace_t* trace) { void FpuUnit::on_tick() { bool idle = true; - auto cur_cycle = SimPlatform::instance().cycles(); for (uint32_t b = 0; b < VX_CFG_NUM_FPU_BLOCKS; ++b) { - auto& inflight = inflight_.at(b); - for (uint32_t i = 0; i < inflight.size();) { - if (inflight.at(i) <= cur_cycle) { - inflight.at(i) = inflight.back(); - inflight.pop_back(); - } else { - ++i; - } - } auto& input = Inputs.at(b); if (!input.empty()) { auto& output = Outputs.at(b); - // A full tag queue stalls admission even though the arithmetic - // pipelines are deeper. - if (!output.full() && inflight.size() < VX_CFG_FPU_QUEUE_SIZE) { + // A full response port (every tag in use) stalls admission even though + // the arithmetic pipelines are deeper. + if (!output.full()) { auto trace = input.peek(); this->execute(trace); - uint32_t delay = this->latency_of(trace); + // Without a lane-gather stage commit reads this port directly, so the + // result also takes the registered stage other units cross into commit. + uint32_t delay = this->latency_of(trace) + (kGather ? 0 : 1); output.send(trace, delay); - inflight.push_back(cur_cycle + delay); input.pop(); } } diff --git a/sim/simx/fpu_unit.h b/sim/simx/fpu_unit.h index 3f5110b485..4de1f44f72 100644 --- a/sim/simx/fpu_unit.h +++ b/sim/simx/fpu_unit.h @@ -13,8 +13,6 @@ #pragma once -#include - #include "func_unit.h" namespace vortex { @@ -23,6 +21,11 @@ class FpuUnit : public FuncUnit { public: FpuUnit(const SimContext& ctx, const char* name, Core*); + // Partial bandwidth: a lane-gather stage takes results off the response + // port. Otherwise the commit arbiter reads the response port directly. + static constexpr bool kGather = (VX_CFG_NUM_FPU_BLOCKS != VX_CFG_ISSUE_WIDTH) + || (VX_CFG_NUM_FPU_LANES != VX_CFG_SIMD_WIDTH); + protected: void on_tick() override; @@ -31,13 +34,6 @@ class FpuUnit : public FuncUnit { void execute(instr_trace_t* trace); uint32_t latency_of(const instr_trace_t* trace) const; - - // Result-exit cycles of operations inside the arithmetic pipelines: an - // operation holds a tag slot from acceptance until its result leaves the - // datapath, so at most VX_CFG_FPU_QUEUE_SIZE operations overlap however - // deep the pipelines are. Results exit out of order across the different - // pipelines, so slots free by exit time, not acceptance order. - std::array, VX_CFG_NUM_FPU_BLOCKS> inflight_; }; } diff --git a/sim/simx/types.h b/sim/simx/types.h index a0133d53e1..1f892713b5 100644 --- a/sim/simx/types.h +++ b/sim/simx/types.h @@ -914,11 +914,6 @@ class IArbiterImpl { IArbiterImpl() {} virtual ~IArbiterImpl() {} virtual uint32_t grant(const BitVector<>& requests) = 0; - virtual uint32_t grant(const BitVector<>& requests, const BitVector<>& suppress) { - // Default: ignore suppress mask, subclasses may override. - __unused (suppress); - return this->grant(requests); - } virtual void reset() = 0; }; @@ -1032,23 +1027,17 @@ class GTOArbiter : public IArbiterImpl { } uint32_t grant(const BitVector<>& requests) override { - BitVector<> no_suppress(size_); - return this->grant(requests, no_suppress); - } - - uint32_t grant(const BitVector<>& requests, const BitVector<>& suppress) override { assert(requests.size() == size_); - assert(suppress.size() == size_); - // greedy: keep granting same requester if still active and unsuppressed - if (last_grant_ < size_ && requests.test(last_grant_) && !suppress.test(last_grant_)) { + // greedy: keep granting same requester if still active + if (last_grant_ < size_ && requests.test(last_grant_)) { this->update_ages(requests, last_grant_); return last_grant_; } - // Then-Oldest: find the unsuppressed requester with the highest age + // Then-Oldest: find the requester with the highest age uint32_t best = -1u; uint32_t best_age = 0; for (uint32_t i = 0; i < size_; ++i) { - if (requests.test(i) && !suppress.test(i) && (best == -1u || age_[i] > best_age)) { + if (requests.test(i) && (best == -1u || age_[i] > best_age)) { best = i; best_age = age_[i]; } @@ -1108,10 +1097,6 @@ class Arbiter { return impl_->grant(requests); } - uint32_t grant(const BitVector<>& requests, const BitVector<>& suppress) { - return impl_->grant(requests, suppress); - } - void reset() { impl_->reset(); } @@ -1589,6 +1574,7 @@ class TxCrossBar : public SimObject> { TxCrossBar( const SimContext& ctx, const char* name, + ArbiterType type, uint32_t num_inputs, uint32_t num_outputs, std::function output_sel, @@ -1597,6 +1583,8 @@ class TxCrossBar : public SimObject> { : SimObject>(ctx, name) , Inputs(num_inputs, this) , Outputs(num_outputs, this) + , type_(type) + , last_grant_(num_outputs, num_inputs - 1) , delay_(delay) , lg2_inputs_(log2ceil(num_inputs)) , lg2_outputs_(log2ceil(num_outputs)) @@ -1606,6 +1594,7 @@ class TxCrossBar : public SimObject> { assert(num_outputs <= 64); assert(ispow2(num_outputs)); assert(output_sel != nullptr); + assert(type == ArbiterType::Priority || type == ArbiterType::RoundRobin); // bypass mode if (num_inputs == 1 && num_outputs == 1) { @@ -1616,10 +1605,11 @@ class TxCrossBar : public SimObject> { TxCrossBar( const SimContext& ctx, const char* name, + ArbiterType type, uint32_t num_inputs, uint32_t num_outputs, std::function output_sel - ) : TxCrossBar(ctx, name, num_inputs, num_outputs, + ) : TxCrossBar(ctx, name, type, num_inputs, num_outputs, output_sel, ((num_inputs > 2) || (num_outputs > 2)) ? 1 : 0) {} @@ -1629,10 +1619,12 @@ class TxCrossBar : public SimObject> { protected: void on_reset() { - //-- + std::fill(last_grant_.begin(), last_grant_.end(), Inputs.size() - 1); } void on_tick(); + ArbiterType type_; + std::vector last_grant_; // per output uint32_t delay_; uint32_t lg2_inputs_; uint32_t lg2_outputs_; @@ -1651,11 +1643,14 @@ void TxCrossBar::on_tick() { return; } - // process incoming requests + // process incoming requests: each output arbitrates on its own, and a + // round-robin grant advances only when the request is accepted. for (uint32_t o = 0; o < O; ++o) { + uint32_t start = (type_ == ArbiterType::RoundRobin) ? (last_grant_.at(o) + 1) : 0; int32_t input_idx = -1; bool has_collision = false; - for (uint32_t i = 0; i < I; ++i) { + for (uint32_t k = 0; k < I; ++k) { + uint32_t i = (start + k) % I; auto& req_in = Inputs.at(i); if (req_in.empty()) continue; @@ -1680,6 +1675,7 @@ void TxCrossBar::on_tick() { if (Outputs.at(o).try_send(RspType(req, input_idx), delay_)) { DT(4, this->name() << " req" << input_idx << "_" << o << ": " << req); req_in.pop(); + last_grant_.at(o) = input_idx; } collisions_ += has_collision; } @@ -1850,7 +1846,7 @@ class TxRxCrossBar : public SimObject> { , lg2_inputs_(log2ceil(num_inputs)) { if (num_inputs != 1 || num_outputs != 1) { - crossbar_ = ReqXbar::Create(name, num_inputs, num_outputs, output_sel, req_delay); + crossbar_ = ReqXbar::Create(name, type, num_inputs, num_outputs, output_sel, req_delay); for (uint32_t i = 0; i < num_inputs; ++i) { ReqIn.at(i).bind(&crossbar_->Inputs.at(i)); } From 3fa7044569db7c4ece62c69d2000eb0ac6684ca4 Mon Sep 17 00:00:00 2001 From: tinebp Date: Thu, 1 Oct 2026 11:26:41 -0700 Subject: [PATCH 3/8] simx: reserve the LSU pending slot when a read enters the request queue The RTL memory scheduler takes a response-buffer entry when the core request is accepted into its queue and frees it at the last response, so loads waiting in the queue and loads in flight share one pool of LSU_PENDING_SIZE. SimX allocated the pending entry only when the request left req_queue, so up to four more reads could sit in the LSU: traces of parity-sgemv show 12 loads inside the SimX LSU against the RTL's 8, and each load waited correspondingly longer (matrix-row loads 59 cycles from LSU entry vs RTL 50). Reads (loads and AMOs) now count against the pool on entry and turn the reservation into their tag on issue; stores and fences are unaffected. Load latency from LSU entry now tracks the RTL (48.4 vs 50.0 and 34.1 vs 31.5 for the two sgemv streams), and parity-sgemv drops from 9.6% to 0.65%, so its 10% tolerance goes back to the 5% default. Co-Authored-By: Claude Opus 5.5 (1M context) --- ci/testcases/core.yaml | 3 --- sim/simx/lsu_unit.cpp | 18 ++++++++++++++++++ sim/simx/lsu_unit.h | 6 ++++++ 3 files changed, 24 insertions(+), 3 deletions(-) diff --git a/ci/testcases/core.yaml b/ci/testcases/core.yaml index 4a2bf69000..921b855606 100644 --- a/ci/testcases/core.yaml +++ b/ci/testcases/core.yaml @@ -146,9 +146,6 @@ tests: # SGEMV: FPU + streaming matrix-vector. - id: parity-sgemv check: model_parity - # Instrs match exactly; cycles diverge ~9.6% (SimX slower). Not yet - # localized. - tolerance: 0.10 # instrs match; cycles ~9.6% via: blackbox app: sgemv args: -m512 -n512 diff --git a/sim/simx/lsu_unit.cpp b/sim/simx/lsu_unit.cpp index 19cb1f0017..78fd3fd466 100644 --- a/sim/simx/lsu_unit.cpp +++ b/sim/simx/lsu_unit.cpp @@ -328,6 +328,13 @@ void LsuUnit::ingest_inputs(uint32_t b) { auto lsu_type_tag = std::get_if(&trace->op_type); if (lsu_type_tag && *lsu_type_tag == LsuType::FENCE && !state.req_queue.empty()) return; + const bool is_read = !lsu_type_tag || *lsu_type_tag == LsuType::LOAD; + if (is_read) { + if (state.pending_reqs.size() + state.staged_reads >= VX_CFG_LSU_PENDING_SIZE) { + return; + } + ++state.staged_reads; + } state.req_queue.push(trace); input.pop(); } @@ -441,6 +448,7 @@ void LsuUnit::process_request_step(uint32_t b) { // drained. if (state.remain_addrs == 0) { this->compute_addrs(b, trace); + state.head_staged = !is_write || is_amo; } // AMO always returns to rd, so it is not direct-commit even though it @@ -540,6 +548,11 @@ void LsuUnit::process_request_step(uint32_t b) { } else { entry_args = std::get(trace->instr_ptr->get_args()); } + if (state.head_staged) { + // the slot reserved at entry becomes this request's tag + --state.staged_reads; + state.head_staged = false; + } tag = state.pending_reqs.allocate({trace, count, is_eop, std::move(lane_entries), entry_args, true}); } lsu_req.tag = tag; @@ -563,6 +576,11 @@ void LsuUnit::process_request_step(uint32_t b) { if (direct_commit) { Outputs.at(b).send(trace); } + if (state.head_staged) { + // a read with no active lane never issues: return its reservation + --state.staged_reads; + state.head_staged = false; + } state.req_queue.pop(); } } diff --git a/sim/simx/lsu_unit.h b/sim/simx/lsu_unit.h index 264e98ebac..02a693998a 100644 --- a/sim/simx/lsu_unit.h +++ b/sim/simx/lsu_unit.h @@ -135,6 +135,10 @@ class LsuUnit : public FuncUnit { // Sized by the outstanding pool (MLP depth), decoupled from the // input staging queue above. HashTable pending_reqs{VX_CFG_LSU_PENDING_SIZE}; + // Reads waiting in req_queue hold a pending_reqs slot from the moment + // they enter it, so staged and outstanding reads share the pool. + uint32_t staged_reads = 0; + bool head_staged = false; FenceController fence; std::vector addr_list; uint32_t remain_addrs = 0; @@ -142,6 +146,8 @@ class LsuUnit : public FuncUnit { void reset() { this->req_queue.clear(); this->pending_reqs.clear(); + this->staged_reads = 0; + this->head_staged = false; this->fence.reset(); this->addr_list.clear(); this->remain_addrs = 0; From 2fc3a479cd20b3577138b439fe96b4c18e2f87f5 Mon Sep 17 00:00:00 2001 From: tinebp Date: Fri, 2 Oct 2026 07:30:16 -0700 Subject: [PATCH 4/8] ci: gate every model_parity case at the 5% default with no known_issue escapes Remove the last known_issue marker (model_parity-wgmma-dxa) and every per-case parity tolerance override (copy-mcast at 10%, the redundant 5% on sgemm-mc), so each model_parity case asserts cycle agreement within the 5% default and a case over the threshold fails CI instead of being reported as an expected failure. The catalog now carries no known_issue anywhere. At the committed SimX model this turns two cases red: wgmma-dxa (SimX 8.5% slower than the RTL) and copy-mcast (6.8%). The pending SimX tile-buffer and LMEM DMA port model brings wgmma-dxa to 2.4%; copy-mcast needs the SimX per-stage latency rework. Co-Authored-By: Claude Opus 5.5 (1M context) --- ci/testcases/core.yaml | 1 - ci/testcases/dxa.yaml | 7 ------- ci/testcases/tensor_wg.yaml | 2 -- 3 files changed, 10 deletions(-) diff --git a/ci/testcases/core.yaml b/ci/testcases/core.yaml index 921b855606..b09ebb949e 100644 --- a/ci/testcases/core.yaml +++ b/ci/testcases/core.yaml @@ -155,7 +155,6 @@ tests: # Tracks to ~3% since SimX models the contended issue-path serializers # (per-fragment load writeback beats, fixed-priority commit arbitration, # and the execution units' bounded in-flight windows). - tolerance: 0.05 via: blackbox app: sgemm args: -n128 diff --git a/ci/testcases/dxa.yaml b/ci/testcases/dxa.yaml index 2941955993..6daca94855 100644 --- a/ci/testcases/dxa.yaml +++ b/ci/testcases/dxa.yaml @@ -459,8 +459,6 @@ tests: threads: [4, 16] - id: model_parity-copy-mcast check: model_parity - # Instrs match; cycles diverge ~6.8% (SimX slower). Not yet localized. - tolerance: 0.10 # instrs match; cycles ~6.8% via: blackbox app: dxa_copy_mcast configs: -DVX_CFG_EXT_DXA_ENABLE @@ -507,11 +505,6 @@ tests: threads: [4, 16] - id: model_parity-wgmma-dxa check: model_parity - # Instructions and DXA transfer timing match; cycles diverge ~8.5% (SimX - # slower), at CTA turnover: SimX drains the C-tile stores later than the RTL - # (store commit point, dispatch queue placement and bank-0 store-miss - # admission). - known_issue: "SimX drains the C-tile stores at CTA turnover slower than the RTL (instrs match; cycles ~8.5%)" via: blackbox app: sgemm_tcu_wg_dxa args: -m 128 -n 128 -k 64 diff --git a/ci/testcases/tensor_wg.yaml b/ci/testcases/tensor_wg.yaml index fd3449444a..4cc78dbf65 100644 --- a/ci/testcases/tensor_wg.yaml +++ b/ci/testcases/tensor_wg.yaml @@ -223,8 +223,6 @@ tests: - id: model_parity-wgmma-fedp2k-rs check: model_parity - # Retired instructions match exactly and the cycle gap is now within tolerance - # (~3.5%): the SimX WGMMA setup/tile-buffer timing closed the former ~8% gap. via: blackbox app: sgemm_tcu_wg args: -m64 -n64 -k64 From 50a5a6f0cf673ade995daa666cc62de656ecf039 Mon Sep 17 00:00:00 2001 From: tinebp Date: Fri, 2 Oct 2026 08:59:31 -0700 Subject: [PATCH 5/8] simx: model the TCU tile buffers and the LMEM DMA port as the RTL builds them wgmma-dxa ran 8.5% slower in SimX than the RTL. The four warps of a warp-group interleaved their C-tile stores one instruction at a time, so the stores evicted each other's lines and data-cache bank 0 fetched every C line two or three times. In the RTL the warps leave the MMA staggered and store back to back. The stagger comes from the TCU's shared B buffer. It holds one bank row, keyed by the lowest block with a WGMMA compute uop, and a block whose (desc_b, step_k, step_n) differs waits; each A buffer holds one k-stripe. SimX planned every A and B line of the WGMMA up front and held all blocks until all were resident, so the blocks ran in lockstep. The tile buffers now refill per key with the RTL's request counts, one read in flight per buffer and fixed-priority arbitration, and a setup uop drops them at the end of its cycle. That alone moved wgmma-dxa only to 8.2%: a refill took 4 cycles in SimX against 1 in the RTL. SimX routed TCU reads and DXA writes through the LSU's per-word bank crossbar, behind the LSU ports and with the TCU ahead of the DXA. Local memory now has the RTL's row-wide DMA port: DXA first, then TCU, one access per cycle, priority over the LSU on the banks it touches, and read data the cycle after the request. SimX vs RTL cycles (model_parity, all 31 cases swept against recorded RTL cycles): wgmma-dxa 8.53% -> 2.43%, wgmma-dxa-mcast 2.26% -> 0.20%, wgmma-sp-dxa 4.21% -> 0.08%, wgmma-int8 0.15% -> 0.32%, wgmma-fedp2k-rs 0.19% -> 0.30%; every other case bit-identical. All 47 functional dxa and tensor_wg SimX cases pass. Co-Authored-By: Claude Opus 5.5 (1M context) --- sim/simx/core.cpp | 26 +- sim/simx/dxa/dxa_core.cpp | 4 +- sim/simx/dxa/dxa_core.h | 4 +- sim/simx/mem/local_mem.cpp | 72 ++++ sim/simx/mem/local_mem.h | 11 + sim/simx/socket.cpp | 10 +- sim/simx/tcu/tcu_tbuf.cpp | 207 ++++++----- sim/simx/tcu/tcu_tbuf.h | 49 +-- sim/simx/tcu/tcu_unit.cpp | 717 +++++++++++++++++++++++-------------- sim/simx/tcu/tcu_unit.h | 6 +- 10 files changed, 714 insertions(+), 392 deletions(-) diff --git a/sim/simx/core.cpp b/sim/simx/core.cpp index 1e447cec6d..37ee5af390 100644 --- a/sim/simx/core.cpp +++ b/sim/simx/core.cpp @@ -122,18 +122,16 @@ class Core::Impl { } // create local memory. - // The DXA drains one full LMEM row per cycle; a row wider than the 64B - // mem_block byteen scope arrives as multiple same-cycle block writes on - // adjacent input ports (see DxaCore::LMEM_PORTS_PER_CORE). + // The DXA writes and the TCU tile-buffer reads share the row-wide DMA + // port, DXA first. snprintf(sname, 100, "%s-lmem", name.c_str()); - uint32_t lmem_num_reqs = LSU_NUM_REQS + VX_CFG_EXT_TCU_ENABLED - + VX_CFG_EXT_DXA_ENABLED * DxaCore::LMEM_PORTS_PER_CORE; local_mem_ = LocalMem::Create(sname, LocalMem::Config{ (1 << VX_CFG_LMEM_LOG_SIZE), LSU_WORD_SIZE, - lmem_num_reqs, + LSU_NUM_REQS, log2ceil(VX_CFG_LMEM_NUM_BANKS), - false + false, + VX_CFG_EXT_DXA_ENABLED + VX_CFG_TCU_WGMMA_ENABLED }); // create lmem switch @@ -254,14 +252,18 @@ class Core::Impl { tcu_unit_ = SimPlatform::instance().create_object(sname, simobject_); func_units_.at((int)FUType::TCU) = tcu_unit_; - // Bind the TCU tile-buffer subsystem (TcuTbuf) to its dedicated LMEM - // port pair, appended after the LSU ports. + #ifdef VX_CFG_TCU_WGMMA_ENABLE + // Bind the TCU tile-buffer subsystem (TcuTbuf) to the LMEM DMA port, + // after the DXA. { auto& tbuf = tcu_unit_->tbuf(); - uint32_t port = LSU_NUM_REQS; - tbuf->lmem_req_out.bind(&local_mem_->Inputs.at(port)); - local_mem_->Outputs.at(port).bind(&tbuf->lmem_rsp_in); + uint32_t base = VX_CFG_EXT_DXA_ENABLED * LocalMem::DMA_PORTS; + for (uint32_t p = 0; p < LocalMem::DMA_PORTS; ++p) { + tbuf->lmem_req_out.at(p).bind(&local_mem_->DmaInputs.at(base + p)); + local_mem_->DmaOutputs.at(base + p).bind(&tbuf->lmem_rsp_in.at(p)); + } } + #endif #ifdef TCU_META_ENABLE // Bind the TCU metadata AGU to the LSU block-0 client port. { diff --git a/sim/simx/dxa/dxa_core.cpp b/sim/simx/dxa/dxa_core.cpp index eb9b42dbb3..5f3c357644 100644 --- a/sim/simx/dxa/dxa_core.cpp +++ b/sim/simx/dxa/dxa_core.cpp @@ -634,7 +634,7 @@ class DxaCore::Impl { // (DxaCore::LMEM_ROW_SIZE bytes — the full banked row the RTL writes in // a single cycle). A row wider than the 64B mem_block byteen scope is // emitted as LMEM_PORTS_PER_CORE same-cycle block writes, one per row - // half on its own LocalMem input port. A row narrower than a block + // half on its own LMEM DMA channel. A row narrower than a block // (small NT) bounds the per-beat gather instead. // A banked-row write is committed atomically, byte enables aside, so a row is @@ -705,7 +705,7 @@ class DxaCore::Impl { break; // next block lands in a different row — next cycle's beat } - // Route this row half to its own LocalMem input port. + // Route this row half to its own LMEM DMA channel. uint32_t half = (DxaCore::LMEM_PORTS_PER_CORE > 1) ? uint32_t((dword / kLmemWordSize) & (DxaCore::LMEM_PORTS_PER_CORE - 1)) : 0u; diff --git a/sim/simx/dxa/dxa_core.h b/sim/simx/dxa/dxa_core.h index 108638614d..7f5c1219ea 100644 --- a/sim/simx/dxa/dxa_core.h +++ b/sim/simx/dxa/dxa_core.h @@ -63,7 +63,7 @@ class DxaCore : public SimObject { // full banked row (DXA_LMEM_WORD_SIZE = LMEM_NUM_BANKS * XLEN/8). The // LocalMem model's byteen scope is one VX_CFG_MEM_BLOCK_SIZE block, so a // row beat is emitted as LMEM_PORTS_PER_CORE same-cycle block writes on - // adjacent input ports (one per row half). + // adjacent LMEM DMA channels (one per row half). static constexpr uint32_t LMEM_ROW_SIZE = VX_CFG_LMEM_NUM_BANKS * (VX_CFG_XLEN / 8); static constexpr uint32_t LMEM_PORTS_PER_CORE = @@ -72,7 +72,7 @@ class DxaCore : public SimObject { // Per-core LMEM write ports (size = VX_CFG_SOCKET_SIZE * // LMEM_PORTS_PER_CORE, core-major). The socket binds each core's - // LocalMem::Inputs[port_dxa + p] here. Write-only — no rsp. + // LocalMem::DmaInputs[p] here. Write-only — no rsp. std::vector> lmem_req_out; DxaCore(const SimContext& ctx, const char* name, Socket* socket); diff --git a/sim/simx/mem/local_mem.cpp b/sim/simx/mem/local_mem.cpp index 9401b83c7a..d2b6d60598 100644 --- a/sim/simx/mem/local_mem.cpp +++ b/sim/simx/mem/local_mem.cpp @@ -108,7 +108,74 @@ class LocalMem::Impl { #endif } + // Serve one DMA client's row access this cycle and return the banks it + // occupies. Clients are served at fixed priority, and DMA wins every bank + // it touches over the LSU: a read touches all banks, a write those its + // byte enables cover. + uint64_t dma_tick() { + uint32_t num_banks = (1 << config_.B); + uint32_t lg2_line_size = log2ceil(config_.line_size); + for (uint32_t c = 0; c < config_.dma_clients; ++c) { + bool pending = false; + bool rsp_ready = true; + for (uint32_t p = 0; p < DMA_PORTS; ++p) { + auto& in = simobject_->DmaInputs.at(c * DMA_PORTS + p); + if (in.empty()) { + continue; + } + pending = true; + if (!in.peek().is_write() && simobject_->DmaOutputs.at(c * DMA_PORTS + p).full()) { + rsp_ready = false; + } + } + if (!pending) { + continue; + } + if (!rsp_ready) { + return 0; + } + uint64_t banks = 0; + for (uint32_t p = 0; p < DMA_PORTS; ++p) { + auto& in = simobject_->DmaInputs.at(c * DMA_PORTS + p); + if (in.empty()) { + continue; + } + auto& req = in.peek(); + uint64_t line_addr = to_local_addr(req.addr) & ~uint64_t(VX_CFG_MEM_BLOCK_SIZE - 1); + if (req.is_write()) { + for (uint32_t b = 0; req.data && b < VX_CFG_MEM_BLOCK_SIZE; ++b) { + if (req.byteen & (1ull << b)) { + uint8_t value = (*req.data)[b]; + ram_.write(&value, line_addr + b, 1); + banks |= 1ull << (((line_addr + b) >> lg2_line_size) & (num_banks - 1)); + } + } +#if VX_CFG_EXT_A_ENABLED + amo_unit_.invalidate(to_local_addr(req.addr), req.hart_id); +#endif + ++perf_stats_.writes; + } else { + MemRsp rsp{req.tag, req.hart_id, req.uuid}; + auto rsp_data = make_mem_block(); + ram_.read(rsp_data->data(), line_addr, VX_CFG_MEM_BLOCK_SIZE); + rsp.data = rsp_data; + // The request arrived over a registered channel; the read + // data returns the cycle after it was issued. + simobject_->DmaOutputs.at(c * DMA_PORTS + p).send(rsp, 0); + banks = (num_banks >= 64) ? ~uint64_t(0) : ((uint64_t(1) << num_banks) - 1); + ++perf_stats_.reads; + } + DT(4, simobject_->name() << "-dma" << c << " req : " << req); + in.pop(); + } + return banks; + } + return 0; + } + void tick() { + uint64_t dma_banks = this->dma_tick(); + // process bank requets from xbar uint32_t num_banks = (1 << config_.B); for (uint32_t i = 0; i < num_banks; ++i) { @@ -126,6 +193,9 @@ class LocalMem::Impl { const uint64_t rdw_addr = amo_rdw_addr_[i]; amo_rdw_valid_[i] = false; #endif + if ((dma_banks >> i) & 1) { + continue; + } auto& xbar_req_out = mem_xbar_->ReqOut.at(i); if (xbar_req_out.empty()) continue; @@ -266,6 +336,8 @@ LocalMem::LocalMem(const SimContext& ctx, const char* name, const Config& config : SimObject(ctx, name) , Inputs(config.num_reqs, this) , Outputs(config.num_reqs, this) + , DmaInputs(config.dma_clients * DMA_PORTS, this) + , DmaOutputs(config.dma_clients * DMA_PORTS, this) , impl_(new Impl(this, config)) {} diff --git a/sim/simx/mem/local_mem.h b/sim/simx/mem/local_mem.h index ada6a2e127..2a879974bd 100644 --- a/sim/simx/mem/local_mem.h +++ b/sim/simx/mem/local_mem.h @@ -19,12 +19,19 @@ namespace vortex { class LocalMem : public SimObject { public: + // A DMA access covers one full bank row. A row wider than a mem_block + // arrives as same-cycle block accesses on DMA_PORTS adjacent channels. + static constexpr uint32_t DMA_ROW_SIZE = VX_CFG_LMEM_NUM_BANKS * (VX_CFG_XLEN / 8); + static constexpr uint32_t DMA_PORTS = + (DMA_ROW_SIZE > VX_CFG_MEM_BLOCK_SIZE) ? (DMA_ROW_SIZE / VX_CFG_MEM_BLOCK_SIZE) : 1; + struct Config { uint32_t capacity; uint32_t line_size; uint32_t num_reqs; uint32_t B; // log2 number of banks bool write_reponse; + uint32_t dma_clients; // row-wide DMA masters, in priority order }; struct PerfStats { @@ -43,6 +50,10 @@ class LocalMem : public SimObject { std::vector> Inputs; std::vector> Outputs; + // DMA port: client c uses channels [c * DMA_PORTS, (c + 1) * DMA_PORTS). + std::vector> DmaInputs; + std::vector> DmaOutputs; + LocalMem(const SimContext& ctx, const char* name, const Config& config); virtual ~LocalMem(); diff --git a/sim/simx/socket.cpp b/sim/simx/socket.cpp index b79aadacf9..73f8bd0625 100644 --- a/sim/simx/socket.cpp +++ b/sim/simx/socket.cpp @@ -323,18 +323,16 @@ class Socket::Impl { sfu->dxa_req_out.bind(&dxa_core_->dxa_req_in.at(c)); } - // DxaCore::lmem_req_out[c] → core's LocalMem.Inputs[port_dxa]. + // DxaCore::lmem_req_out[c] → core's LocalMem DMA port (first client). // A tx_callback on the channel fires barrier_event_release for each // DXA-write packet carrying notify_done at the cycle LMEM receives it. - uint32_t port_dxa = LSU_NUM_REQS; - #ifdef VX_CFG_EXT_TCU_ENABLE - port_dxa += 1; - #endif + static_assert(DxaCore::LMEM_PORTS_PER_CORE == LocalMem::DMA_PORTS, + "a DXA row write spans the LMEM DMA ports"); for (uint32_t c = 0; c < cores_per_socket; ++c) { Core* core = cores_.at(c).get(); for (uint32_t p = 0; p < DxaCore::LMEM_PORTS_PER_CORE; ++p) { auto& ch = dxa_core_->lmem_req_out.at(c * DxaCore::LMEM_PORTS_PER_CORE + p); - ch.bind(&core->local_mem()->Inputs.at(port_dxa + p)); + ch.bind(&core->local_mem()->DmaInputs.at(p)); ch.tx_callback([core](const MemReq& req, uint64_t /*cycles*/) { if (req.is_write() && req.flags.dxa_notify_done) { // notify_bar_id arrives in raw (encoded) form: low byte = cta_no, diff --git a/sim/simx/tcu/tcu_tbuf.cpp b/sim/simx/tcu/tcu_tbuf.cpp index 4e9d794eaa..e80594297d 100644 --- a/sim/simx/tcu/tcu_tbuf.cpp +++ b/sim/simx/tcu/tcu_tbuf.cpp @@ -15,7 +15,6 @@ #include "constants.h" #include "debug.h" #include -#include #include #include @@ -25,58 +24,58 @@ namespace { constexpr uint64_t kLineMask = ~uint64_t(VX_CFG_MEM_BLOCK_SIZE - 1); -// Q+1 sources fan into one external LMEM port. Source IDs: +// Q+1 buffers share one LMEM port. Buffer IDs, in priority order: // 0 .. VX_CFG_NUM_TCU_BLOCKS-1 → abuf[b] // VX_CFG_NUM_TCU_BLOCKS → bbuf -constexpr uint32_t kNumSources = VX_CFG_NUM_TCU_BLOCKS + 1; +constexpr uint32_t kNumBuffers = VX_CFG_NUM_TCU_BLOCKS + 1; constexpr uint32_t kAOffset = 0; constexpr uint32_t kBOffset = VX_CFG_NUM_TCU_BLOCKS; -// Per-source line cache. Resident, in-flight and pending state are tracked -// independently per source; the wrapper arbitrates the shared LMEM port. -struct LineBuf { - std::deque pending_q_; - std::unordered_map inflight_; // per-source tag → addr - std::unordered_map> resident_; +struct Buffer { + std::deque pending_; + std::unordered_map inflight_; // tag → line + std::unordered_map> lines_; + uint64_t fill_cycle_ = 0; + bool rsp_seen_ = false; // a response arrived this cycle uint32_t next_tag_ = 0; uint64_t reads_ = 0; - void plan(const std::vector& line_addrs) { - std::unordered_set inflight_set; - for (auto& kv : inflight_) inflight_set.insert(kv.second); - for (auto a : line_addrs) { - uint64_t line = a & kLineMask; - if (resident_.count(line)) continue; - if (inflight_set.count(line)) continue; - pending_q_.push_back(line); - inflight_set.insert(line); - } + bool filling() const { + return !pending_.empty() || !inflight_.empty(); } - bool ready() const { - return pending_q_.empty() && inflight_.empty(); + void fill(std::vector&& reads) { + this->invalidate(); + for (auto& r : reads) { + pending_.push_back(std::move(r)); + } + fill_cycle_ = SimPlatform::instance().cycles(); } std::shared_ptr read(uint64_t line_addr) const { - auto it = resident_.find(line_addr & kLineMask); - if (it == resident_.end()) return nullptr; + auto it = lines_.find(line_addr & kLineMask); + if (it == lines_.end()) { + return nullptr; + } return it->second; } + // A response still in flight for an abandoned fill finds no tag and is dropped. void invalidate() { - resident_.clear(); + pending_.clear(); + inflight_.clear(); + lines_.clear(); } void reset() { - pending_q_.clear(); - inflight_.clear(); - resident_.clear(); + this->invalidate(); + rsp_seen_ = false; next_tag_ = 0; reads_ = 0; } }; -// Pack/unpack the source ID alongside the per-source tag in MemReq::tag. +// Pack the buffer ID alongside the per-buffer tag in MemReq::tag. constexpr uint32_t kSrcShift = 16; constexpr uint32_t kSubTagMask = (1u << kSrcShift) - 1; @@ -93,83 +92,119 @@ class TcuTbuf::Impl { Impl(TcuTbuf* simobject) : simobject_(simobject) {} void reset() { - for (auto& b : bufs_) b.reset(); - rr_next_ = 0; - } - - void plan(uint32_t source, const std::vector& line_addrs) { - bufs_.at(source).plan(line_addrs); - } - - bool ready(uint32_t source) const { - return bufs_.at(source).ready(); + for (auto& b : bufs_) { + b.reset(); + } } - std::shared_ptr read(uint32_t source, uint64_t line_addr) const { - return bufs_.at(source).read(line_addr); + Buffer& buf(uint32_t source) { + return bufs_.at(source); } - void invalidate(uint32_t source) { - bufs_.at(source).invalidate(); + const Buffer& buf(uint32_t source) const { + return bufs_.at(source); } uint64_t reads() const { uint64_t total = 0; - for (auto& b : bufs_) total += b.reads_; + for (auto& b : bufs_) { + total += b.reads_; + } return total; } void tick() { - // 1) drain one response and route it to the source that issued it. - auto& rsp = simobject_->lmem_rsp_in; - if (!rsp.empty()) { + for (auto& b : bufs_) { + b.rsp_seen_ = false; + } + + for (auto& rsp : simobject_->lmem_rsp_in) { + if (rsp.empty()) { + continue; + } auto& r = rsp.peek(); uint32_t source = unpack_source(r.tag); - uint32_t sub_tag = unpack_sub_tag(r.tag); - if (source < kNumSources) { - auto& buf = bufs_.at(source); - auto it = buf.inflight_.find(sub_tag); - if (it != buf.inflight_.end()) { - if (r.data) buf.resident_[it->second] = r.data; - buf.inflight_.erase(it); + if (source < kNumBuffers) { + auto& b = bufs_.at(source); + auto it = b.inflight_.find(unpack_sub_tag(r.tag)); + if (it != b.inflight_.end()) { + if (r.data) { + b.lines_[it->second] = r.data; + } + b.inflight_.erase(it); + b.rsp_seen_ = true; + if (!b.filling()) { + DT(3, simobject_->name() << " " << this->buf_name(source) << ": READY"); + } } } rsp.pop(); } - // 2) round-robin pick one source with pending work and submit one req. - auto& req = simobject_->lmem_req_out; - if (req.full()) return; - for (uint32_t i = 0; i < kNumSources; ++i) { - uint32_t s = (rr_next_ + i) % kNumSources; - auto& buf = bufs_.at(s); - if (buf.pending_q_.empty()) continue; - uint64_t addr = buf.pending_q_.front(); - // inflight_ is keyed by the tag as it appears on the wire, so the - // counter must be masked here and not only inside pack_tag(). - uint32_t sub_tag = (buf.next_tag_++) & kSubTagMask; - uint32_t tag = pack_tag(s, sub_tag); - MemReq m(MemOp::LD, addr, /*data*/nullptr, /*byteen*/0, tag, /*hart_id*/0, /*uuid*/0); - m.flags.local = 1; // TCU TBUF reads from LMEM - req.send(m, 1); - buf.inflight_[sub_tag] = addr; - buf.pending_q_.pop_front(); - ++buf.reads_; - rr_next_ = (s + 1) % kNumSources; + // An abuf issues its next read in the cycle its previous one returns; + // the bbuf issues it the cycle after. + uint64_t cycle = SimPlatform::instance().cycles(); + for (uint32_t s = 0; s < kNumBuffers; ++s) { + auto& b = bufs_.at(s); + if (b.pending_.empty() || !b.inflight_.empty()) { + continue; + } + if (cycle <= b.fill_cycle_) { + continue; + } + if (s == kBOffset && b.rsp_seen_) { + continue; + } + this->issue(s, b); break; } } private: + void issue(uint32_t source, Buffer& b) { + auto& lines = b.pending_.front(); + uint32_t n = std::min(lines.size(), LMEM_PORTS); + for (uint32_t p = 0; p < n; ++p) { + if (simobject_->lmem_req_out.at(p).full()) { + return; + } + } + for (uint32_t p = 0; p < n; ++p) { + uint64_t addr = lines.at(p); + // inflight_ is keyed by the tag as it appears on the wire, so the + // counter must be masked here and not only inside pack_tag(). + uint32_t sub_tag = (b.next_tag_++) & kSubTagMask; + MemReq m(MemOp::LD, addr, /*data*/nullptr, /*byteen*/0, + pack_tag(source, sub_tag), /*hart_id*/0, /*uuid*/0); + m.flags.local = 1; + simobject_->lmem_req_out.at(p).send(m, 1); + b.inflight_[sub_tag] = addr; + } + DT(3, simobject_->name() << " " << this->buf_name(source) + << ": rd_req addr=0x" << std::hex << lines.front() << std::dec + << ", lines=" << n); + // A read wider than the port set finishes as a further read. + if (n < lines.size()) { + lines.erase(lines.begin(), lines.begin() + n); + } else { + b.pending_.pop_front(); + } + ++b.reads_; + } + + std::string buf_name(uint32_t source) const { + return (source == kBOffset) ? std::string("bbuf") + : ("abuf" + std::to_string(source - kAOffset)); + } + TcuTbuf* simobject_; - std::array bufs_; - uint32_t rr_next_ = 0; + std::array bufs_; }; TcuTbuf::TcuTbuf(const SimContext& ctx, const char* name) : SimObject(ctx, name) - , lmem_req_out(this) - , lmem_rsp_in(this) + , lmem_req_out(LMEM_PORTS, this) + , lmem_rsp_in(LMEM_PORTS, this) , impl_(new Impl(this)) {} @@ -178,24 +213,24 @@ TcuTbuf::~TcuTbuf() { delete impl_; } void TcuTbuf::on_reset() { impl_->reset(); } void TcuTbuf::on_tick() { impl_->tick(); } -void TcuTbuf::plan_a(uint32_t b, const std::vector& line_addrs) { - impl_->plan(kAOffset + b, line_addrs); +void TcuTbuf::fill_a(uint32_t b, std::vector reads) { + impl_->buf(kAOffset + b).fill(std::move(reads)); } -void TcuTbuf::plan_b(const std::vector& line_addrs) { - impl_->plan(kBOffset, line_addrs); +void TcuTbuf::fill_b(std::vector reads) { + impl_->buf(kBOffset).fill(std::move(reads)); } -bool TcuTbuf::ready_a(uint32_t b) const { return impl_->ready(kAOffset + b); } -bool TcuTbuf::ready_b() const { return impl_->ready(kBOffset); } +bool TcuTbuf::filling_a(uint32_t b) const { return impl_->buf(kAOffset + b).filling(); } +bool TcuTbuf::filling_b() const { return impl_->buf(kBOffset).filling(); } std::shared_ptr TcuTbuf::read_a(uint32_t b, uint64_t line_addr) const { - return impl_->read(kAOffset + b, line_addr); + return impl_->buf(kAOffset + b).read(line_addr); } std::shared_ptr TcuTbuf::read_b(uint64_t line_addr) const { - return impl_->read(kBOffset, line_addr); + return impl_->buf(kBOffset).read(line_addr); } -void TcuTbuf::invalidate_a(uint32_t b) { impl_->invalidate(kAOffset + b); } -void TcuTbuf::invalidate_b() { impl_->invalidate(kBOffset); } +void TcuTbuf::invalidate_a(uint32_t b) { impl_->buf(kAOffset + b).invalidate(); } +void TcuTbuf::invalidate_b() { impl_->buf(kBOffset).invalidate(); } uint64_t TcuTbuf::reads() const { return impl_->reads(); } diff --git a/sim/simx/tcu/tcu_tbuf.h b/sim/simx/tcu/tcu_tbuf.h index 96f1bced1d..91d3f32224 100644 --- a/sim/simx/tcu/tcu_tbuf.h +++ b/sim/simx/tcu/tcu_tbuf.h @@ -14,50 +14,59 @@ #pragma once #include "types.h" +#include "local_mem.h" namespace vortex { // TCU tile-buffer subsystem. // -// Owns the A / B operand line caches for WGMMA and arbitrates their LMEM -// traffic onto a single LMEM port pair. +// Holds the WGMMA operand storage and fetches it from LMEM: // -// abuf × Q per-block A-tile line cache (one per warp lane) -// bbuf × 1 shared B-tile line cache (TB-shared across all Q blocks) +// abuf × Q per-block A k-stripe +// bbuf × 1 B bank row, shared by all Q blocks // -// Sparse metadata is preloaded into the TcuUnit's per-warp `sparse_meta_` -// SRAM via TCU_LD ahead of the MMA dispatch; it does not flow through this buffer. +// The consumer (TcuUnit) owns the refill keys: it decides when a buffer +// refills and hands over the refill as a list of LMEM reads. Each read is +// one bank-row request on the hardware port and covers one or more +// mem_block lines here, sent together on parallel ports. A buffer keeps one +// read in flight; buffers share the port at fixed priority, abuf[0] first +// and bbuf last. // -// All line caches are address-keyed. The consumer (TcuUnit) plans the set -// of line addresses each cache should fetch for the current WGMMA, waits -// for the corresponding `ready_*()` to clear, then extracts operand bytes -// from the resident `mem_block_t` payloads via `read_*()`. +// Sparse metadata is preloaded into the TcuUnit's per-warp `sparse_meta_` +// SRAM via TCU_LD ahead of the MMA dispatch; it does not flow through here. class TcuTbuf : public SimObject { public: using Ptr = std::shared_ptr; - // Single external LMEM port pair. Internal arbitration fans Q+1 sources - // (abuf × Q, bbuf × 1) into one bank-row request channel. - SimChannel lmem_req_out; - SimChannel lmem_rsp_in; + // A bank row wider than a mem_block is read as same-cycle block reads on + // adjacent LMEM DMA channels. + static constexpr uint32_t LMEM_PORTS = LocalMem::DMA_PORTS; + + // Lines of one bank-row read. + using LmemRead = std::vector; + + std::vector> lmem_req_out; + std::vector> lmem_rsp_in; TcuTbuf(const SimContext& ctx, const char* name); virtual ~TcuTbuf(); - // Plan/query/read API per role. `b` indexes the per-block lane. - void plan_a(uint32_t b, const std::vector& line_addrs); - void plan_b(const std::vector& line_addrs); + // Drop the buffer's contents and fetch `reads` in order. `b` indexes the + // per-block A buffer. + void fill_a(uint32_t b, std::vector reads); + void fill_b(std::vector reads); - bool ready_a(uint32_t b) const; - bool ready_b() const; + bool filling_a(uint32_t b) const; + bool filling_b() const; std::shared_ptr read_a(uint32_t b, uint64_t line_addr) const; std::shared_ptr read_b(uint64_t line_addr) const; + // Drop the buffer's contents and abandon a fill in progress. void invalidate_a(uint32_t b); void invalidate_b(); - // Total LMEM port-cycles issued since last reset (perf counter). + // Bank-row reads issued since the last reset (perf counter). uint64_t reads() const; protected: diff --git a/sim/simx/tcu/tcu_unit.cpp b/sim/simx/tcu/tcu_unit.cpp index 398aa6ed42..e31b65ac68 100644 --- a/sim/simx/tcu/tcu_unit.cpp +++ b/sim/simx/tcu/tcu_unit.cpp @@ -46,6 +46,27 @@ static constexpr uint32_t kFedpLatency = VX_CFG_TCU_LATENCY; // unblocked result bypasses the landing queue). static constexpr uint32_t kMmaLatency = 1 + kFedpLatency; +static constexpr uint64_t kLineMask = ~uint64_t(VX_CFG_MEM_BLOCK_SIZE - 1); + +// WGMMA tile-buffer geometry in 32-bit words. +static constexpr uint32_t kLmemBanks = VX_CFG_LMEM_NUM_BANKS; +static constexpr uint32_t kBankRowWords = kLmemBanks * (VX_CFG_XLEN / 32); +static constexpr uint32_t kBBlockWords = kFedpWords * cfg::tcN; +static constexpr uint32_t kBBlockWordsSp = cfg::tcK * cfg::tcN * 2; +// One bank row holds kBSubBlocks consecutive dense B blocks. +static constexpr uint32_t kBSubBlocks = + (VX_CFG_NUM_THREADS > kBBlockWords) ? (VX_CFG_NUM_THREADS / kBBlockWords) : 1; +// A sparse B block spanning two 32-bit bank rows takes two reads. +static constexpr bool kBSparseTwoFetch = + (kBBlockWordsSp == 2 * kLmemBanks) && (VX_CFG_XLEN == 32); +// A block-major A stripe is m_steps blocks, each padded to whole bank rows. +static constexpr uint32_t a_stripe_reads(uint32_t block_words) { + return (wg_cfg::m_steps * ((block_words + kLmemBanks - 1) / kLmemBanks) * kLmemBanks + + kBankRowWords - 1) / kBankRowWords; +} +static constexpr uint32_t kAStripeReads = a_stripe_reads(cfg::tcM * kFedpWords); +static constexpr uint32_t kAStripeReadsSp = a_stripe_reads(cfg::tcM * cfg::tcK); + inline uint64_t nan_box(uint32_t value) { return value | 0xffffffff00000000; } @@ -352,7 +373,6 @@ class TcuUnit::Impl { , perf_stats_() { exec_done_.fill(false); - wgmma_planned_warps_.fill(0); in_wgmma_.fill(false); wgmma_desc_.fill({0, 0}); } @@ -373,12 +393,14 @@ class TcuUnit::Impl { } #endif exec_done_.fill(false); - wgmma_planned_warps_.fill(0); in_wgmma_.fill(false); lmem_desc_.clear(); wgmma_desc_.fill({0, 0}); cta_owner_a_.fill(-1); - cta_owner_b_ = -1; + #ifdef VX_CFG_TCU_WGMMA_ENABLE + abuf_.fill(abuf_slot_t{}); + bbuf_ = bbuf_slot_t{}; + #endif cur_block_ = 0; #ifdef TCU_META_ENABLE agu_.fill(agu_state_t{}); @@ -534,104 +556,8 @@ class TcuUnit::Impl { this->agu_step(); #endif #ifdef VX_CFG_TCU_WGMMA_ENABLE - // Q-warp lock-step probe. - // Pass 1 — identify active WGMMA blocks and prime each one's plan() on - // first uop. WMMA and TCU_LD blocks are unaffected (no Q-coupling). - uint32_t wgmma_active = 0; - for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { - auto& input = simobject_->Inputs.at(b); - if (input.empty()) continue; - auto trace = input.peek(); - if (!tcu_is_wgmma(std::get(trace->op_type))) continue; - - uint32_t wid = trace->wid; - uint64_t wid_bit = (uint64_t(1) << wid); - int32_t new_cta = (int32_t)core_->scheduler().warp(wid).cta_csrs.cta_id; - auto& instr = *trace->instr_ptr; - auto tpuArgs = std::get(instr.get_args()); - - if (tpuArgs.is_setup_uop) - continue; - - wgmma_active |= (1u << b); - - // CTA-overlap fence — defer this block's WGMMA if any other block - // is mid-flight with a different CTA. The shared B buffer assumes - // single-CTA occupancy across all blocks. - bool block_other_cta_inflight = false; - for (uint32_t k = 0; k < VX_CFG_NUM_TCU_BLOCKS; ++k) { - if (k == b) continue; - if (in_wgmma_.at(k) && cta_owner_a_.at(k) != new_cta) { - block_other_cta_inflight = true; - break; - } - } - if (block_other_cta_inflight) { - wgmma_active &= ~(1u << b); - continue; - } - - if (wgmma_planned_warps_.at(b) & wid_bit) continue; - if (!(tpuArgs.step_m == 0 && tpuArgs.step_n == 0 && tpuArgs.step_k == 0)) { - // Non-first uop arrived without a prior plan: first uop already drained. - // Mark planned and continue (descriptors persist in lmem_desc_[wid]). - wgmma_planned_warps_.at(b) |= wid_bit; - continue; - } - uint32_t a_desc = wgmma_desc_[wid][0]; - uint32_t b_desc = wgmma_desc_[wid][1]; - bool needs_setup = kFedp2K && !tpuArgs.is_a_smem - && (std::get(trace->op_type) != TcuType::WGMMA_SP); - if (tpuArgs.is_first_uop && !needs_setup) { - if (tpuArgs.is_a_smem) - a_desc = trace->src_data.at(0).at(0).u32; - b_desc = trace->src_data.at(1).at(0).u32; - wgmma_desc_[wid][0] = a_desc; - wgmma_desc_[wid][1] = b_desc; - } - - // Drop the shared B buffer only when no other block is mid-WGMMA — - // otherwise we'd evict their resident bytes mid-flight. - bool any_in_wgmma = false; - for (auto v : in_wgmma_) any_in_wgmma = any_in_wgmma || v; - auto& tbuf = simobject_->tbuf(); - if (!any_in_wgmma) { - tbuf->invalidate_b(); - cta_owner_b_ = -1; - } - // Only drop the per-block A buffer when no warp is currently in flight. - if (!in_wgmma_.at(b)) { - tbuf->invalidate_a(b); - } - this->plan_wgmma_lines(b, wid, a_desc, b_desc, tpuArgs, - std::get(trace->op_type) == TcuType::WGMMA_SP); - if (tbuf->ready_a(b) && tbuf->ready_b()) { - ++perf_stats_.tbuf_cache_hits; - } - in_wgmma_.at(b) = true; - wgmma_planned_warps_.at(b) |= wid_bit; - cta_owner_a_.at(b) = new_cta; - if (cta_owner_b_ == -1) cta_owner_b_ = new_cta; - } - - // Pass 2 — all active WGMMA blocks must have A/B operands resident - // before any of them advances. - if (wgmma_active != 0) { - uint32_t ready_mask = 0; - auto& tbuf = simobject_->tbuf(); - for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { - if (!((wgmma_active >> b) & 1u)) continue; - auto trace = simobject_->Inputs.at(b).peek(); - auto tpuArgs = std::get(trace->instr_ptr->get_args()); - bool a_ok = !tpuArgs.is_a_smem || tbuf->ready_a(b); - bool b_ok = tbuf->ready_b(); - if (a_ok && b_ok) ready_mask |= (1u << b); - } - if (ready_mask != wgmma_active) { - ++perf_stats_.tbuf_stalls; - return; // hold all blocks; per-block dispatch deferred to next tick - } - } + uint32_t wgmma_ready = this->tbuf_step(); + uint32_t setup_fired = 0; #endif for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { @@ -642,23 +568,11 @@ class TcuUnit::Impl { auto tcu_type = std::get(trace->op_type); auto tpuArgs = std::get(trace->instr_ptr->get_args()); - #ifdef VX_CFG_TCU_WGMMA_ENABLE - // CTA-overlap fence deferred this block — skip until pass 1 plans it. - if (tcu_is_wgmma(tcu_type) && !tpuArgs.is_setup_uop && - !(wgmma_planned_warps_.at(b) & (uint64_t(1) << trace->wid))) + #ifdef VX_CFG_TCU_WGMMA_ENABLE + // A WGMMA compute uop fires once its tile-buffer operands are ready. + if (tcu_is_wgmma(tcu_type) && !tpuArgs.is_setup_uop + && !exec_done_.at(b) && !((wgmma_ready >> b) & 1u)) { continue; - if (tcu_is_wgmma(tcu_type) && !tpuArgs.is_setup_uop) { - int32_t this_cta = (int32_t)core_->scheduler().warp(trace->wid).cta_csrs.cta_id; - bool block_other_cta_inflight = false; - for (uint32_t k = 0; k < VX_CFG_NUM_TCU_BLOCKS; ++k) { - if (k == b) continue; - if (in_wgmma_.at(k) && cta_owner_a_.at(k) != this_cta) { - block_other_cta_inflight = true; - break; - } - } - if (block_other_cta_inflight) - continue; } #endif @@ -687,10 +601,10 @@ class TcuUnit::Impl { uint32_t a_desc = rs1_data.empty() ? 0 : rs1_data.at(0).u32; uint32_t b_desc = rs2_data.empty() ? 0 : rs2_data.at(0).u32; cur_block_ = b; + int32_t this_cta = (int32_t)core_->scheduler().warp(wid).cta_csrs.cta_id; // CTA lockstep invariant: no block may execute a WGMMA uop for a // different cta_id while another block is mid-WGMMA. if (!tpuArgs.is_setup_uop) { - int32_t this_cta = (int32_t)core_->scheduler().warp(wid).cta_csrs.cta_id; for (uint32_t k = 0; k < VX_CFG_NUM_TCU_BLOCKS; ++k) { if (k == b) continue; if (in_wgmma_.at(k) && cta_owner_a_.at(k) != this_cta) { @@ -707,6 +621,19 @@ class TcuUnit::Impl { a_desc, b_desc, rs1_data, rs2_data, rs3_data, rd_data, tcu_is_sparse(tcu_type), tpuArgs.cd_nregs, tpuArgs.is_a_smem, tpuArgs.is_setup_uop); + if (tpuArgs.is_setup_uop) { + setup_fired |= (1u << b); + } else { + // The block owns its CTA from its first compute uop to its last. + if (tpuArgs.is_first_uop) { + in_wgmma_.at(b) = true; + cta_owner_a_.at(b) = this_cta; + } + if (tpuArgs.is_last_uop) { + in_wgmma_.at(b) = false; + cta_owner_a_.at(b) = -1; + } + } } break; #endif #ifdef TCU_META_ENABLE @@ -756,18 +683,6 @@ class TcuUnit::Impl { #endif if (simobject_->Outputs.at(b).try_send(trace, delay)) { exec_done_.at(b) = false; - #ifdef VX_CFG_TCU_WGMMA_ENABLE - // Clear this warp's plan bit on its last uop so the next WGMMA - // re-decodes descriptors. Block stays in_wgmma_ until all warps drain. - if (tcu_is_wgmma(tcu_type) && trace->instr_ptr->get_fu_unlock()) { - uint64_t wid_bit = (uint64_t(1) << trace->wid); - wgmma_planned_warps_.at(b) &= ~wid_bit; - if (wgmma_planned_warps_.at(b) == 0) { - in_wgmma_.at(b) = false; - cta_owner_a_.at(b) = -1; - } - } - #endif #ifdef TCU_META_ENABLE // TCU_LD retired: free the AGU for the next metadata load. if (tcu_type == TcuType::TCU_LD) { @@ -778,97 +693,366 @@ class TcuUnit::Impl { input.pop(); } } + + #ifdef VX_CFG_TCU_WGMMA_ENABLE + // A setup uop drops its block's A stripe and the shared B row at the end + // of the cycle; uops firing alongside it still read the old contents. + if (setup_fired != 0) { + auto& tbuf = simobject_->tbuf(); + for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + if ((setup_fired >> b) & 1u) { + abuf_.at(b) = abuf_slot_t{}; + tbuf->invalidate_a(b); + } + } + bbuf_ = bbuf_slot_t{}; + tbuf->invalidate_b(); + } + #endif } - // Plan all line addresses required for the current WGMMA's A, B and - // sparse-metadata tiles into the per-role caches inside TcuTbuf. - // Lines already resident or in-flight are skipped (additive plan). - void plan_wgmma_lines(uint32_t b, uint32_t wid, - uint32_t a_desc, uint32_t b_desc, - const IntrTcuArgs& args, bool is_sparse) { - uint32_t fmt_s = args.fmt_s; - bool is_a_smem = args.is_a_smem; - uint32_t e_bits = elem_bits(fmt_s); - // NRC: cd_nregs 0/1/2 → 8/16/32; xtileN = NRC * NT / xtileM. - uint32_t nrc = (args.cd_nregs == 0) ? 8 : (args.cd_nregs == 1) ? 16 : 32; - uint32_t xtile_n = (nrc * VX_CFG_NUM_THREADS) / wg_cfg::xtileM; +#ifdef VX_CFG_TCU_WGMMA_ENABLE + // A block's A buffer: one k-stripe. + struct abuf_slot_t { + bool valid = false; + bool fetching = false; + bool first_refetched = false; + uint32_t step_k = 0; + }; - lmem_desc_t sd_a{}, sd_b{}; - if (is_a_smem) { - sd_a = {uint64_t(VX_MEM_LMEM_BASE_ADDR) + (a_desc & 0xFFFF), (a_desc >> 16) * 8 / e_bits, false}; - lmem_desc_[wid][0] = sd_a; + // The shared B buffer and its refill key. + struct bbuf_slot_t { + bool valid = false; + bool fetching = false; + bool first_refetched = false; + bool kmajor = false; + bool sparse = false; + uint64_t row = 0; // block-major: LMEM address of the bank row + int32_t cta = -1; // K-major: the (cta, step_k, step_n) block + uint32_t step_k = 0; + uint32_t step_n = 0; + + bool same_row(const bbuf_slot_t& key) const { + if (kmajor != key.kmajor || sparse != key.sparse) { + return false; + } + if (kmajor) { + return cta == key.cta && step_k == key.step_k && step_n == key.step_n; + } + return row == key.row; } - sd_b = {uint64_t(VX_MEM_LMEM_BASE_ADDR) + (b_desc & 0xFFFF), (b_desc >> 16) * 8 / e_bits, false}; - lmem_desc_[wid][1] = sd_b; - - // tileK = xtileK × ratio (ratio = 32/e_bits); sparse compresses K on A only. - uint32_t ratio = 32 / e_bits; - uint32_t tile_k = uint32_t(wg_cfg::xtileK) * ratio; - uint32_t a_k = is_sparse ? (tile_k / 2) : tile_k; - - // ldm==0 → block-major layout; ldm!=0 → row-major (stride in elements). - uint32_t fedp_words = kFedpWords; - uint32_t b_k_blk_dim = fedp_words * ratio; - uint32_t a_k_blk_dim = is_sparse ? (cfg::tcK * ratio) : b_k_blk_dim; - uint32_t a_blk_elems = cfg::tcM * a_k_blk_dim; - uint32_t b_blk_elems = b_k_blk_dim * cfg::tcN; - uint32_t n_steps = xtile_n / cfg::tcN; + }; + // One block's WGMMA compute uop as the tile buffers see it. + struct tbuf_req_t { + bool valid = false; + uint32_t wid = 0; + int32_t cta = -1; + bool is_first_uop = false; + bool is_sparse = false; + bool a_is_smem = false; + uint32_t step_m = 0; + uint32_t step_n = 0; + uint32_t step_k = 0; + uint32_t cd_nregs = 0; + uint32_t fmt_s = 0; + uint32_t desc_b = 0; + + // The A buffer forces a refetch on a WGMMA's first compute uop. + bool first_compute() const { + return step_m == 0 && step_n == 0 && step_k == 0; + } + }; + + // Evaluate the A/B tile buffers for the uops at the head of each block + // and return the blocks whose WGMMA compute uop may fire. + // + // The B buffer holds one bank row and refills when the key of the lowest + // block with a uop changes; a block whose (desc_b, step_k, step_n) differs + // from that block's waits. Each A buffer holds its block's k-stripe. + // Refills start from the buffer state at the beginning of the cycle, and a + // refill that completes this cycle serves from the next. + uint32_t tbuf_step() { auto& tbuf = simobject_->tbuf(); - // Plan A lines (SS mode only): xtileM rows × a_k columns. + std::array reqs; + int32_t rep = -1; + for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + auto& input = simobject_->Inputs.at(b); + if (input.empty() || exec_done_.at(b)) { + continue; + } + auto trace = input.peek(); + auto tcu_type = std::get(trace->op_type); + if (!tcu_is_wgmma(tcu_type)) { + continue; + } + auto tpuArgs = std::get(trace->instr_ptr->get_args()); + if (tpuArgs.is_setup_uop) { + continue; + } + uint32_t wid = trace->wid; + int32_t cta = (int32_t)core_->scheduler().warp(wid).cta_csrs.cta_id; + // CTA lockstep: wait while another block is mid-WGMMA for another CTA. + bool cta_conflict = false; + for (uint32_t k = 0; k < VX_CFG_NUM_TCU_BLOCKS; ++k) { + if (k != b && in_wgmma_.at(k) && cta_owner_a_.at(k) != cta) { + cta_conflict = true; + break; + } + } + if (cta_conflict) { + continue; + } + + bool is_sparse = (tcu_type == TcuType::WGMMA_SP); + bool needs_setup = kFedp2K && !tpuArgs.is_a_smem && !is_sparse; + if (tpuArgs.is_first_uop && !needs_setup) { + // The first compute uop carries the live descriptors. + if (tpuArgs.is_a_smem) { + wgmma_desc_[wid][0] = trace->src_data.at(0).at(0).u32; + } + wgmma_desc_[wid][1] = trace->src_data.at(1).at(0).u32; + this->decode_descs(wid, tpuArgs.fmt_s, tpuArgs.is_a_smem); + } + + auto& req = reqs.at(b); + req.valid = true; + req.wid = wid; + req.cta = cta; + req.is_first_uop = tpuArgs.is_first_uop; + req.is_sparse = is_sparse; + req.a_is_smem = tpuArgs.is_a_smem; + req.step_m = tpuArgs.step_m; + req.step_n = tpuArgs.step_n; + req.step_k = tpuArgs.step_k; + req.cd_nregs = tpuArgs.cd_nregs; + req.fmt_s = tpuArgs.fmt_s; + req.desc_b = wgmma_desc_[wid][1]; + if (rep < 0) { + rep = b; + } + } + + // Shared B buffer, keyed by the lowest block with a uop. + bool bbuf_ready = true; + if (rep >= 0) { + auto& r = reqs.at(rep); + bbuf_slot_t key = this->bbuf_key(r); + bool resident = bbuf_.valid && bbuf_.same_row(key) + && (!key.kmajor || !r.is_first_uop || bbuf_.first_refetched); + if (resident) { + ++perf_stats_.tbuf_cache_hits; + } else { + bbuf_ready = false; + ++perf_stats_.tbuf_stalls; + if (!bbuf_.fetching) { + bool first_refetched = bbuf_.first_refetched; + bbuf_ = key; + bbuf_.fetching = true; + bbuf_.first_refetched = first_refetched; + DT(3, simobject_->name() << " bbuf: alloc desc_b=0x" << std::hex << r.desc_b + << std::dec << ", sparse=" << r.is_sparse << ", step_k=" << r.step_k + << ", step_n=" << r.step_n << ", row=0x" << std::hex << key.row + << std::dec << ", wid=" << r.wid); + tbuf->fill_b(this->b_refill(r)); + } + } + } + + // Per-block A buffers. + std::array abuf_ready; + for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + auto& r = reqs.at(b); + auto& slot = abuf_.at(b); + abuf_ready.at(b) = true; + if (!r.valid || !r.a_is_smem) { + continue; + } + bool resident = slot.valid && slot.step_k == r.step_k + && (!r.first_compute() || slot.first_refetched); + if (resident) { + continue; + } + abuf_ready.at(b) = false; + ++perf_stats_.tbuf_stalls; + if (!slot.fetching) { + slot.valid = false; + slot.fetching = true; + slot.step_k = r.step_k; + tbuf->fill_a(b, this->a_refill(r)); + } + } + + uint32_t ready_mask = 0; + for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + auto& r = reqs.at(b); + if (!r.valid) { + continue; + } + auto& p = reqs.at(rep); + bool key_match = ((r.desc_b & 0xffff) == (p.desc_b & 0xffff)) + && r.step_k == p.step_k + && r.step_n == p.step_n + && r.cd_nregs == p.cd_nregs; + if (abuf_ready.at(b) && bbuf_ready && key_match) { + ready_mask |= (1u << b); + } + } + + // Buffer state for the next cycle. + if (bbuf_.fetching && !tbuf->filling_b()) { + bbuf_.fetching = false; + bbuf_.valid = true; + if (rep >= 0 && reqs.at(rep).is_first_uop) { + bbuf_.first_refetched = true; + } + } else if (rep >= 0 && bbuf_ready && !reqs.at(rep).is_first_uop) { + bbuf_.first_refetched = false; + } + for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + auto& r = reqs.at(b); + auto& slot = abuf_.at(b); + if (slot.fetching && !tbuf->filling_a(b)) { + slot.fetching = false; + slot.valid = true; + if (r.valid && r.first_compute()) { + slot.first_refetched = true; + } + } else if (r.valid && abuf_ready.at(b) && !r.first_compute()) { + slot.first_refetched = false; + } + } + + return ready_mask; + } + + // Decode a warp's WGMMA descriptors into LMEM base and row stride. + void decode_descs(uint32_t wid, uint32_t fmt_s, bool is_a_smem) { + uint32_t e_bits = elem_bits(fmt_s); + uint32_t a_desc = wgmma_desc_[wid][0]; + uint32_t b_desc = wgmma_desc_[wid][1]; if (is_a_smem) { - bool a_block_major = (sd_a.ldm == 0); - std::vector a_lines; - a_lines.reserve(uint32_t(wg_cfg::xtileM) * a_k); - for (uint32_t r = 0; r < wg_cfg::xtileM; ++r) { - for (uint32_t c = 0; c < a_k; ++c) { - uint64_t elem_off; - if (a_block_major) { - uint32_t m_blk = r / cfg::tcM; - uint32_t i_in = r % cfg::tcM; - uint32_t k_blk = c / a_k_blk_dim; - uint32_t k_in = c % a_k_blk_dim; - elem_off = (k_blk * wg_cfg::m_steps + m_blk) * a_blk_elems - + i_in * a_k_blk_dim + k_in; - } else { - elem_off = uint64_t(r) * sd_a.ldm + c; + lmem_desc_[wid][0] = {uint64_t(VX_MEM_LMEM_BASE_ADDR) + (a_desc & 0xFFFF), (a_desc >> 16) * 8 / e_bits, false}; + } + lmem_desc_[wid][1] = {uint64_t(VX_MEM_LMEM_BASE_ADDR) + (b_desc & 0xFFFF), (b_desc >> 16) * 8 / e_bits, false}; + } + + static uint32_t xtile_n_of(uint32_t cd_nregs) { + // NRC: cd_nregs 0/1/2 → 8/16/32; xtileN = NRC * NT / xtileM. + uint32_t nrc = (cd_nregs == 0) ? 8 : (cd_nregs == 1) ? 16 : 32; + return (nrc * VX_CFG_NUM_THREADS) / wg_cfg::xtileM; + } + + // Refill key of the B row serving `req`. Block-major B is keyed by the + // bank row holding the (step_k, step_n) block; K-major B by the block. + bbuf_slot_t bbuf_key(const tbuf_req_t& req) { + const auto& sd_b = lmem_desc_[req.wid][1]; + bbuf_slot_t key; + key.kmajor = (sd_b.ldm != 0); + key.sparse = req.is_sparse; + if (key.kmajor) { + key.cta = req.cta; + key.step_k = req.step_k; + key.step_n = req.step_n; + } else { + uint32_t n_steps = xtile_n_of(req.cd_nregs) / cfg::tcN; + uint32_t blk = req.step_k * n_steps + req.step_n; + uint32_t blk_bytes = req.is_sparse ? (kBBlockWordsSp * 4) : (kBBlockWords * 4); + uint32_t blks_per_row = req.is_sparse ? 1 : kBSubBlocks; + key.row = sd_b.base + uint64_t(blk / blks_per_row) * blks_per_row * blk_bytes; + } + return key; + } + + // LMEM reads that refill the B buffer for `req`: one bank row for a + // block-major row (two when a sparse block spans two rows) and one per + // N-row of the block in K-major layout. + std::vector b_refill(const tbuf_req_t& req) { + const auto& sd_b = lmem_desc_[req.wid][1]; + uint32_t e_bits = elem_bits(req.fmt_s); + uint32_t ratio = 32 / e_bits; + uint32_t xtile_n = xtile_n_of(req.cd_nregs); + uint32_t n_steps = xtile_n / cfg::tcN; + uint32_t k_words = req.is_sparse ? cfg::tcK : kFedpWords; + bool kmajor = (sd_b.ldm != 0); + + // (step_k, step_n) blocks this refill brings in. + uint32_t blk = req.step_k * n_steps + req.step_n; + uint32_t blks = (kmajor || req.is_sparse) ? 1 : kBSubBlocks; + uint32_t first_blk = blk - (blk % blks); + + std::vector lines; + for (uint32_t bi = first_blk; bi < first_blk + blks; ++bi) { + uint32_t step_k = bi / n_steps; + uint32_t step_n = bi % n_steps; + for (uint32_t j = 0; j < cfg::tcN; ++j) { + uint32_t col = step_n * cfg::tcN + j; + for (uint32_t z = 0; z < k_words; ++z) { + uint32_t k_elem = (step_k * k_words + z) * ratio * (req.is_sparse ? 2 : 1); + uint32_t k_elems = req.is_sparse ? (2 * ratio) : ratio; + for (uint32_t e = 0; e < k_elems; ++e) { + uint64_t off = elem_offset(sd_b, k_elem + e, col, e_bits, xtile_n, + true, req.is_sparse, false); + lines.push_back((sd_b.base + off * e_bits / 8) & kLineMask); } - uint64_t addr = sd_a.base + elem_off * e_bits / 8; - a_lines.push_back(addr & ~uint64_t(VX_CFG_MEM_BLOCK_SIZE - 1)); } } - tbuf->plan_a(b, a_lines); - // Sparse metadata is preloaded into sparse_meta_ via TCU_LD; - // no metadata lines are planned through tbuf here. } - // Plan B lines: always dense in K, tileK rows × xtileN columns. - // ldm == 0 → block-major; ldm != 0 → K-major (smem[n*ldm + k]). - bool b_block_major = (sd_b.ldm == 0); - std::vector b_lines; - b_lines.reserve(tile_k * xtile_n); - for (uint32_t r = 0; r < tile_k; ++r) { - for (uint32_t c = 0; c < xtile_n; ++c) { - uint64_t elem_off; - if (b_block_major) { - uint32_t k_blk = r / b_k_blk_dim; - uint32_t r_in = r % b_k_blk_dim; - uint32_t n_blk = c / cfg::tcN; - uint32_t n_in = c % cfg::tcN; - // Within-block layout: N outer, K inner. - elem_off = (k_blk * n_steps + n_blk) * b_blk_elems - + n_in * b_k_blk_dim + r_in; - } else { - elem_off = uint64_t(c) * sd_b.ldm + r; + uint32_t num_reads = kmajor ? cfg::tcN + : (req.is_sparse && kBSparseTwoFetch) ? 2 : 1; + return split_reads(std::move(lines), num_reads); + } + + // LMEM reads that refill a block's A buffer with the k-stripe of `req`: + // one per bank row of a block-major stripe, one per row in row-major. + std::vector a_refill(const tbuf_req_t& req) { + const auto& sd_a = lmem_desc_[req.wid][0]; + uint32_t e_bits = elem_bits(req.fmt_s); + uint32_t ratio = 32 / e_bits; + uint32_t xtile_n = xtile_n_of(req.cd_nregs); + uint32_t k_words = req.is_sparse ? cfg::tcK : kFedpWords; + + std::vector lines; + for (uint32_t row = 0; row < wg_cfg::m_steps * cfg::tcM; ++row) { + for (uint32_t z = 0; z < k_words; ++z) { + uint32_t k_elem = (req.step_k * k_words + z) * ratio; + for (uint32_t e = 0; e < ratio; ++e) { + uint64_t off = elem_offset(sd_a, row, k_elem + e, e_bits, xtile_n, + false, false, req.is_sparse); + lines.push_back((sd_a.base + off * e_bits / 8) & kLineMask); } - uint64_t addr = sd_b.base + elem_off * e_bits / 8; - b_lines.push_back(addr & ~uint64_t(VX_CFG_MEM_BLOCK_SIZE - 1)); } } - tbuf->plan_b(b_lines); + + uint32_t num_reads = (sd_a.ldm != 0) ? (wg_cfg::m_steps * cfg::tcM) + : req.is_sparse ? kAStripeReadsSp : kAStripeReads; + return split_reads(std::move(lines), num_reads); } + // Spread a refill's distinct lines over the hardware's read count. + static std::vector split_reads(std::vector lines, + uint32_t num_reads) { + std::sort(lines.begin(), lines.end()); + lines.erase(std::unique(lines.begin(), lines.end()), lines.end()); + uint32_t n = lines.size(); + if (n == 0) { + return {}; + } + std::vector reads(num_reads); + for (uint32_t i = 0; i < num_reads; ++i) { + if (n >= num_reads) { + reads.at(i).assign(lines.begin() + (uint64_t(i) * n) / num_reads, + lines.begin() + (uint64_t(i + 1) * n) / num_reads); + } else { + reads.at(i).push_back(lines.at(i % n)); + } + } + return reads; + } +#endif // VX_CFG_TCU_WGMMA_ENABLE + void wmma(uint32_t wid, uint32_t fmt_s, @@ -1185,6 +1369,68 @@ class TcuUnit::Impl { #endif } + // Element offset of one operand element from its descriptor base. A is + // (row=M, col=K) and B is (row=K, col=N); `pack_along_row` selects B. + static uint64_t elem_offset(const lmem_desc_t& desc, uint32_t row, uint32_t col, + uint32_t e_bits, uint32_t xtile_n, + bool pack_along_row, bool sparse_b, + bool sparse_a_layout) { + uint32_t ratio = (e_bits >= 32) ? 1 : (32 / e_bits); + uint64_t elem_off; + if (desc.ldm == 0) { + // Block-major SMEM. K dimension is along col for A (pack_along_row + // false) and along row for B (pack_along_row true). + uint32_t k_blk_dim = (sparse_a_layout && !pack_along_row) + ? (cfg::tcK * ratio) + : (kFedpWords * ratio); + if (pack_along_row && sparse_b) { + // Sparse B in flat (candidate-pair) layout: block-contiguous + // K-word-major / N-inner order [kw_in*tcN + n_in]. + uint32_t b_tcK_words = cfg::tcK * 2; + uint32_t k_word = row / ratio; + uint32_t elem = row % ratio; + uint32_t k_blk = k_word / b_tcK_words; + uint32_t kw_in = k_word % b_tcK_words; + uint32_t n_blk = col / cfg::tcN; + uint32_t n_in = col % cfg::tcN; + uint32_t blk_words = cfg::tcN * b_tcK_words; + uint32_t n_steps = xtile_n / cfg::tcN; + uint64_t word_off = (k_blk * n_steps + n_blk) * blk_words + + (kw_in * cfg::tcN + n_in); + elem_off = word_off * ratio + elem; + } else if (pack_along_row) { + // Dense B (block-major): r is K coord, c is N coord; N outer, K inner. + uint32_t k_blk = row / k_blk_dim; + uint32_t r_in = row % k_blk_dim; + uint32_t n_blk = col / cfg::tcN; + uint32_t n_in = col % cfg::tcN; + uint32_t b_blk_elems = k_blk_dim * cfg::tcN; + uint32_t n_steps = xtile_n / cfg::tcN; + elem_off = (k_blk * n_steps + n_blk) * b_blk_elems + + n_in * k_blk_dim + r_in; + } else { + // A: r is M coord, c is K coord. + uint32_t m_blk = row / cfg::tcM; + uint32_t i_in = row % cfg::tcM; + uint32_t k_blk = col / k_blk_dim; + uint32_t k_in = col % k_blk_dim; + uint32_t a_blk_elems = cfg::tcM * k_blk_dim; + elem_off = (k_blk * wg_cfg::m_steps + m_blk) * a_blk_elems + + i_in * k_blk_dim + k_in; + } + } else if (desc.col_major) { + elem_off = uint64_t(col) * desc.ldm + row; + } else if (pack_along_row) { + // B: K-major (N-outer K-inner). row=K, col=N; + // ldm = stride in elements between N rows. + elem_off = uint64_t(col) * desc.ldm + row; + } else { + // A: row-major (M-outer K-inner). row=M, col=K. + elem_off = uint64_t(row) * desc.ldm + col; + } + return elem_off; + } + // Gather one 32-bit operand word from a TCU line cache. // `read_line` is supplied by the caller and routes to the right per-role // buffer inside TcuTbuf (A → read_a, B → read_b). For sub-32-bit formats, @@ -1201,60 +1447,8 @@ class TcuUnit::Impl { for (uint32_t r = 0; r < ratio; ++r) { uint32_t cur_row = pack_along_row ? (row + r) : row; uint32_t cur_col = pack_along_row ? col : (col + r); - uint64_t elem_off; - if (desc.ldm == 0) { - // Block-major SMEM. K dimension is along col for A (pack_along_row - // false) and along row for B (pack_along_row true). - uint32_t k_blk_dim = (sparse_a_layout && !pack_along_row) - ? (cfg::tcK * ratio) - : (kFedpWords * ratio); - if (pack_along_row && sparse_b) { - // Sparse B in flat (candidate-pair) layout: block-contiguous - // K-word-major / N-inner order [kw_in*tcN + n_in] (mirrors - // vx_tensor.h b_sp_flat_idx). The bbuf applies the candidate-pair - // read-perm; here we just read logically. - uint32_t b_tcK_words = cfg::tcK * 2; - uint32_t k_word = cur_row / ratio; - uint32_t elem = cur_row % ratio; - uint32_t k_blk = k_word / b_tcK_words; - uint32_t kw_in = k_word % b_tcK_words; - uint32_t n_blk = cur_col / cfg::tcN; - uint32_t n_in = cur_col % cfg::tcN; - uint32_t blk_words = cfg::tcN * b_tcK_words; - uint32_t n_steps = cur_xtile_n_ / cfg::tcN; - uint64_t word_off = (k_blk * n_steps + n_blk) * blk_words - + (kw_in * cfg::tcN + n_in); - elem_off = word_off * ratio + elem; - } else if (pack_along_row) { - // Dense B (block-major): r is K coord, c is N coord; N outer, K inner. - uint32_t k_blk = cur_row / k_blk_dim; - uint32_t r_in = cur_row % k_blk_dim; - uint32_t n_blk = cur_col / cfg::tcN; - uint32_t n_in = cur_col % cfg::tcN; - uint32_t b_blk_elems = k_blk_dim * cfg::tcN; - uint32_t n_steps = cur_xtile_n_ / cfg::tcN; - elem_off = (k_blk * n_steps + n_blk) * b_blk_elems - + n_in * k_blk_dim + r_in; - } else { - // A: r is M coord, c is K coord. - uint32_t m_blk = cur_row / cfg::tcM; - uint32_t i_in = cur_row % cfg::tcM; - uint32_t k_blk = cur_col / k_blk_dim; - uint32_t k_in = cur_col % k_blk_dim; - uint32_t a_blk_elems = cfg::tcM * k_blk_dim; - elem_off = (k_blk * wg_cfg::m_steps + m_blk) * a_blk_elems - + i_in * k_blk_dim + k_in; - } - } else if (desc.col_major) { - elem_off = uint64_t(cur_col) * desc.ldm + cur_row; - } else if (pack_along_row) { - // B: K-major (N-outer K-inner). cur_row=K, cur_col=N; - // ldm = stride in elements between N rows. - elem_off = uint64_t(cur_col) * desc.ldm + cur_row; - } else { - // A: row-major (M-outer K-inner). cur_row=M, cur_col=K. - elem_off = uint64_t(cur_row) * desc.ldm + cur_col; - } + uint64_t elem_off = elem_offset(desc, cur_row, cur_col, e_bits, cur_xtile_n_, + pack_along_row, sparse_b, sparse_a_layout); uint64_t byte_addr = desc.base + elem_off * e_bits / 8; auto line = read_line(byte_addr); if (!line) { @@ -1355,8 +1549,6 @@ class TcuUnit::Impl { mutable PerfStats perf_stats_; // Per-block guard: execute already happened for this trace; reset on pop(). std::array exec_done_; - // Per-block bitmask of warp IDs with planned WGMMA lines; cleared on fu_unlock. - std::array wgmma_planned_warps_; // True while a block is between its first and last WGMMA uop. std::array in_wgmma_; // Current block index, set before delegating to wgmma(). @@ -1366,9 +1558,12 @@ class TcuUnit::Impl { bool cur_is_sparse_ = false; // xtileN for the active WGMMA (derived from NRC). uint32_t cur_xtile_n_ = 8; - // CTA owner per block's A buffer and the shared B buffer (-1 = unowned). + // CTA each block holds while mid-WGMMA (-1 = none). std::array cta_owner_a_{}; - int32_t cta_owner_b_ = -1; +#ifdef VX_CFG_TCU_WGMMA_ENABLE + std::array abuf_; + bbuf_slot_t bbuf_; +#endif }; /////////////////////////////////////////////////////////////////////////////// diff --git a/sim/simx/tcu/tcu_unit.h b/sim/simx/tcu/tcu_unit.h index fc78a33e06..93328060f4 100644 --- a/sim/simx/tcu/tcu_unit.h +++ b/sim/simx/tcu/tcu_unit.h @@ -60,9 +60,9 @@ class TcuUnit : public FuncUnit { struct PerfStats { uint64_t latency = 0; - uint64_t tbuf_stalls = 0; // cycles stalled on TcuTbufA/TcuSharedB readiness - uint64_t tbuf_cache_hits = 0; // WGMMA entries with all lines already resident (cross-WGMMA reuse) - uint64_t lmem_reads = 0; // sum of TcuTbufA + TcuSharedB LmemReqs issued + uint64_t tbuf_stalls = 0; // cycles a WGMMA uop waits on its A stripe or the B row, summed over buffers + uint64_t tbuf_cache_hits = 0; // cycles the resident B row serves the representative block + uint64_t lmem_reads = 0; // bank-row reads issued by the A and B buffers PerfStats& operator+=(const PerfStats& rhs) { this->latency += rhs.latency; From 34962a44662b3bf3e90c4996d0c43b42ee6a9904 Mon Sep 17 00:00:00 2001 From: tinebp Date: Fri, 2 Oct 2026 09:00:07 -0700 Subject: [PATCH 6/8] simx: release stalled warps and wake dependents on the RTL's schedule copy-mcast ran 6.8% slower in SimX. Measured per opcode, every gap was in the ops that hold their warp until they resolve: a branch took 18 cycles from schedule to the warp's next schedule against the RTL's 17, SPLIT and WSYNC 20 against 17, JOIN 20 against 18. The RTL releases such a warp from the unit that resolves it, through one register (a join through one more, for the divergence-stack pop); SimX released it only when the result left the functional unit. The ALU and SFU now send the release to the core on a one-cycle channel, and a TMC that disables every thread retires its warp the same way. Releasing on time exposed three places where SimX ran faster than the RTL, compared per event on vecadd: - a warp freed by TMC was reused for the next CTA in the same cycle; the RTL picks from a registered active set, registers the pick, and the new warp is schedulable the cycle after; - integer multiply/divide had the plain ALU's latency; the multiplier pipeline is two stages longer; - a dependent instruction dispatched 4 cycles after its producer committed, against 7 in the RTL, while ALU and SFU results reached commit 2 cycles late. The scoreboard now clears a destination two cycles after commit, and the ALU and SFU result latencies lose the matching two cycles. SimX vs RTL cycles, all 31 model_parity cases (on top of the previous commit): copy-mcast 6.83% -> 2.22%, softmax 3.44% -> 0.04%, copy 4.06% -> 0.35%, sgemm2-dxa 2.83% -> 0.81%, vecadd 2.94% -> 1.54%; 30 of 31 within 5%. Known regression: parity-sgemm-mc goes from 2.24% to 11.85% (SimX slower) and fails the 5% gate. Most of it follows the scoreboard release delay (5.11% without it). On a reduced run FMADD dispatch-to-commit averages 49 cycles in SimX against 15 in the RTL: SimX dispatches the FMA chain early and it queues in front of the FPU and at commit, where the RTL holds it in the instruction buffer. Under investigation; raycast (4.53%), draw3d (4.73%) and sparse-fp16 (4.89%) are close to the limit. Co-Authored-By: Claude Opus 5.5 (1M context) --- sim/simx/alu_unit.cpp | 26 ++++++++++++++++++-------- sim/simx/alu_unit.h | 3 +++ sim/simx/core.cpp | 23 +++++++++++++++++++++++ sim/simx/core.h | 5 +++++ sim/simx/instr_trace.h | 2 ++ sim/simx/scheduler.cpp | 25 ++++++++++++++++++++----- sim/simx/scheduler.h | 6 ++++++ sim/simx/scoreboard.cpp | 19 +++++++++++++++++-- sim/simx/scoreboard.h | 14 ++++++++++++++ sim/simx/sfu_unit.cpp | 15 +++++++++++++-- sim/simx/sfu_unit.h | 4 ++++ sim/simx/types.h | 7 +++++++ sim/simx/wctl_unit.cpp | 10 ++++++++-- sim/simx/wctl_unit.h | 8 +++++--- 14 files changed, 145 insertions(+), 22 deletions(-) diff --git a/sim/simx/alu_unit.cpp b/sim/simx/alu_unit.cpp index 648729c0ad..4da3a189fd 100644 --- a/sim/simx/alu_unit.cpp +++ b/sim/simx/alu_unit.cpp @@ -26,8 +26,12 @@ using namespace vortex; AluUnit::AluUnit(const SimContext& ctx, const char* name, Core* core) : FuncUnit(ctx, name, core) + , branch_ctl_out(this, VX_CFG_NUM_WARPS) {} +// Cycles from this unit's execute to its output, beyond the commit-queue hop. +// Execute stands for the integer ALU's registered result (the branch_ctl +// stage), so a plain op leaves at once. uint32_t AluUnit::latency_of(const instr_trace_t* trace) const { if (std::get_if(&trace->op_type)) { auto alu_type = std::get(trace->op_type); @@ -45,16 +49,16 @@ uint32_t AluUnit::latency_of(const instr_trace_t* trace) const { case AluType::AND: case AluType::OR: case AluType::CZERO: - return 2; + return 0; default: std::abort(); } } else if (std::get_if(&trace->op_type)) { - return 2; + return 0; } else if (std::get_if(&trace->op_type)) { - return 2; + return 0; } else if (std::get_if(&trace->op_type)) { - return 2; + return 0; } else if (std::get_if(&trace->op_type)) { auto br_type = std::get(trace->op_type); switch (br_type) { @@ -62,24 +66,24 @@ uint32_t AluUnit::latency_of(const instr_trace_t* trace) const { case BrType::JAL: case BrType::JALR: case BrType::SYS: - return 2; + return 0; default: std::abort(); } } else if (std::get_if(&trace->op_type)) { auto mdv_type = std::get(trace->op_type); switch (mdv_type) { + // The multiplier pipeline is three stages deep against the integer + // ALU's single response stage; simulation divides run in that same + // pipeline rather than iteratively. case MdvType::MUL: case MdvType::MULHU: case MdvType::MULH: case MdvType::MULHSU: - return 2; case MdvType::DIV: case MdvType::DIVU: case MdvType::REM: case MdvType::REMU: - // Simulation divides are fully pipelined at the multiplier's - // depth, not iterative. return 2; default: std::abort(); @@ -562,6 +566,12 @@ void AluUnit::on_tick() { if (!output.full()) { auto trace = input.peek(); this->execute(trace); + // A branch releases its warp as it resolves, not when its result + // retires. + if (std::get_if(&trace->op_type) && trace->eop && trace->resume_warp) { + branch_ctl_out.send(trace->wid, 1); + trace->resume_warp = false; + } uint32_t delay = this->latency_of(trace); output.send(trace, delay); input.pop(); diff --git a/sim/simx/alu_unit.h b/sim/simx/alu_unit.h index b5f5288fba..b94904cc97 100644 --- a/sim/simx/alu_unit.h +++ b/sim/simx/alu_unit.h @@ -21,6 +21,9 @@ class AluUnit : public FuncUnit { public: AluUnit(const SimContext& ctx, const char* name, Core*); + // Resolved branch of a stalled warp, registered toward the scheduler. + SimChannel branch_ctl_out; + protected: void on_tick() override; diff --git a/sim/simx/core.cpp b/sim/simx/core.cpp index 37ee5af390..cdaf880714 100644 --- a/sim/simx/core.cpp +++ b/sim/simx/core.cpp @@ -247,6 +247,8 @@ class Core::Impl { func_units_.at((int)FUType::LSU) = SimPlatform::instance().create_object(sname, simobject_); snprintf(sname, 100, "%s-sfu", name.c_str()); func_units_.at((int)FUType::SFU) = SimPlatform::instance().create_object(sname, simobject_); + std::static_pointer_cast(func_units_.at((int)FUType::ALU))->branch_ctl_out.bind(&simobject_->branch_ctl_in); + std::static_pointer_cast(func_units_.at((int)FUType::SFU))->warp_ctl_out.bind(&simobject_->warp_ctl_in); #ifdef VX_CFG_EXT_TCU_ENABLE snprintf(sname, 100, "%s-tcu", name.c_str()); tcu_unit_ = SimPlatform::instance().create_object(sname, simobject_); @@ -711,6 +713,25 @@ class Core::Impl { } void commit() { + // A resolved branch or warp-control op releases its warp, or retires it + // when it disabled all its threads, a cycle after it resolves; the + // scheduler observes the change the cycle after that. + auto& branch_ctl = simobject_->branch_ctl_in; + while (!branch_ctl.empty()) { + scheduler_->resume(branch_ctl.peek()); + branch_ctl.pop(); + } + auto& warp_ctl = simobject_->warp_ctl_in; + while (!warp_ctl.empty()) { + auto& ctl = warp_ctl.peek(); + if (ctl.exit) { + scheduler_->setTmask(ctl.wid, ThreadMask(VX_CFG_NUM_THREADS)); + } else { + scheduler_->resume(ctl.wid); + } + warp_ctl.pop(); + } + // Fan-in: route per-block FU outputs to per-iw commit queues by trace->wid. // Each FU has NUM_*_BLOCKS outputs; the original iw was lost during // dispatcher aggregation, so we recover it from the warp id and try_send @@ -1069,6 +1090,8 @@ Core::Core(const SimContext& ctx, , icache_rsp_in(1, this) , dcache_req_out(VX_CFG_DCACHE_NUM_REQS, this) , dcache_rsp_in(VX_CFG_DCACHE_NUM_REQS, this) + , branch_ctl_in(this, VX_CFG_NUM_WARPS) + , warp_ctl_in(this, VX_CFG_NUM_WARPS) , gbar_arrive_out(this) , gbar_resume_in(this) #ifdef VX_CFG_EXT_RASTER_ENABLE diff --git a/sim/simx/core.h b/sim/simx/core.h index 0d346215d7..df1e117b7b 100644 --- a/sim/simx/core.h +++ b/sim/simx/core.h @@ -78,6 +78,11 @@ class Core : public SimObject { std::vector> dcache_req_out; std::vector> dcache_rsp_in; + // Warp releases from resolved branches (ALU) and warp-control ops (SFU). + // A warp has at most one release in flight, so NUM_WARPS bounds each. + SimChannel branch_ctl_in; + SimChannel warp_ctl_in; + // Global-barrier event links, wired core <-> cluster at elaboration. SimEventLink gbar_arrive_out; SimEventLink gbar_resume_in; diff --git a/sim/simx/instr_trace.h b/sim/simx/instr_trace.h index bf15096734..59c808f4f5 100644 --- a/sim/simx/instr_trace.h +++ b/sim/simx/instr_trace.h @@ -67,6 +67,8 @@ struct instr_trace_t { // Set by a func-unit when a fetch_stall instruction has resolved; the warp // is released when the trace drains from the FU output (commit fan-in). + // Branches and warp-control ops release their warp as they resolve instead + // and leave this clear. bool resume_warp; uint64_t issue_time ; diff --git a/sim/simx/scheduler.cpp b/sim/simx/scheduler.cpp index 4ccf32ff25..2f6acac77a 100644 --- a/sim/simx/scheduler.cpp +++ b/sim/simx/scheduler.cpp @@ -90,6 +90,8 @@ void Scheduler::on_reset() { stalled_warps_.reset(); stalled_warps_next_.reset(); active_warps_.reset(); + cta_active_view_.reset(); + cta_fire_ = false; // Sequencers live on Core now; Core::on_reset() resets them. wspawn_.valid = false; @@ -158,12 +160,21 @@ void Scheduler::activate_warp(uint32_t wid, const cta_warp_record_t& rec) { instr_trace_t* Scheduler::schedule(const WarpMask& warp_mask) { int scheduled_warp = -1; - // Dispatch one CTA warp + // Dispatch one CTA warp. The dispatcher picks from the registered active + // set and its pick lands the next cycle; the warp it activates becomes + // schedulable the cycle after that. + if (cta_fire_) { + activate_warp(cta_fire_wid_, cta_fire_rec_); + stalled_warps_.set(cta_fire_wid_); + cta_fire_ = false; + } { uint32_t wid; cta_warp_record_t rec; - if (cta_dispatcher_->step(active_warps_, &wid, &rec)) { - activate_warp(wid, rec); + if (cta_dispatcher_->step(cta_active_view_, &wid, &rec)) { + cta_fire_ = true; + cta_fire_wid_ = wid; + cta_fire_rec_ = rec; } } @@ -240,11 +251,15 @@ instr_trace_t* Scheduler::schedule(const WarpMask& warp_mask) { // becomes visible to the pick loop next cycle — so a warp released as its // instruction resolves is never re-scheduled the same cycle. stalled_warps_ = stalled_warps_next_; + cta_active_view_ = active_warps_; + if (cta_fire_) { + cta_active_view_.set(cta_fire_wid_); + } return trace; } bool Scheduler::running() const { - return active_warps_.any() || cta_dispatcher_->running() + return active_warps_.any() || cta_fire_ || cta_dispatcher_->running() #ifdef VX_CFG_EXT_RASTER_ENABLE || fwd_armed_ #endif @@ -404,7 +419,7 @@ void Scheduler::fwd_try_inject() { // Find any free warp slot — there is no driver warp to skip in the push model. int wid = -1; for (uint32_t w = 0; w < VX_CFG_NUM_WARPS; ++w) { - if (!active_warps_.test(w)) { wid = int(w); break; } + if (!active_warps_.test(w) && !(cta_fire_ && w == cta_fire_wid_)) { wid = int(w); break; } } if (wid < 0) break; // no free slot this cycle diff --git a/sim/simx/scheduler.h b/sim/simx/scheduler.h index c02556d333..a397eb60ba 100644 --- a/sim/simx/scheduler.h +++ b/sim/simx/scheduler.h @@ -211,6 +211,12 @@ class Scheduler : public SimObject { WarpMask active_warps_; WarpMask stalled_warps_; // registered (current) state read by schedule() WarpMask stalled_warps_next_; // next-state written by suspend()/resume() + // CTA dispatch pipeline: the active set the dispatcher sees, and the warp + // it picked last cycle. + WarpMask cta_active_view_; + bool cta_fire_ = false; + uint32_t cta_fire_wid_ = 0; + cta_warp_record_t cta_fire_rec_; uint32_t ipdom_size_; wspawn_t wspawn_; uint32_t mpm_class_; diff --git a/sim/simx/scoreboard.cpp b/sim/simx/scoreboard.cpp index 30ec252ef6..fee0fece6d 100644 --- a/sim/simx/scoreboard.cpp +++ b/sim/simx/scoreboard.cpp @@ -35,6 +35,20 @@ void Scoreboard::on_reset() { } owners_.clear(); commit_counts_.clear(); + pending_releases_.clear(); +} + +void Scoreboard::on_tick() { + uint64_t now = SimPlatform::instance().cycles(); + while (!pending_releases_.empty() && pending_releases_.front().due <= now) { + auto& r = pending_releases_.front(); + owners_.erase(get_reg_id(r.reg, r.wid)); + in_use_regs_.at(r.wid).at((int)r.reg.type).reset(r.reg.idx); + pending_releases_.pop_front(); + } + if (pending_releases_.empty()) { + this->tick_sleep(); + } } bool Scoreboard::in_use(instr_trace_t* trace) const { @@ -89,9 +103,10 @@ void Scoreboard::release(instr_trace_t* trace) { assert(trace->wb); assert(in_use_regs_.at(trace->wid).at((int)trace->dst_reg.type).test(trace->dst_reg.idx)); assert(owners_.count(reg_id) != 0); - owners_.erase(reg_id); commit_counts_.erase(reg_id); - in_use_regs_.at(trace->wid).at((int)trace->dst_reg.type).reset(trace->dst_reg.idx); + pending_releases_.push_back({trace->wid, trace->dst_reg, + SimPlatform::instance().cycles() + kReleaseDelay}); + this->tick_wake(); } bool Scoreboard::commit_packet(instr_trace_t* trace) { diff --git a/sim/simx/scoreboard.h b/sim/simx/scoreboard.h index 4eeece71fc..2cad40ea5b 100644 --- a/sim/simx/scoreboard.h +++ b/sim/simx/scoreboard.h @@ -15,6 +15,7 @@ #include "types.h" #include "instr_trace.h" +#include #include #include @@ -40,6 +41,9 @@ class Scoreboard : public SimObject { void reserve(instr_trace_t* trace); + // The destination is written to the register file the cycle after + // commit and clears in the scoreboard the cycle after that; a dependent + // instruction sees it free from then on. void release(instr_trace_t* trace); // Per-packet commit notifier. Returns true when every SIMD-split packet @@ -52,13 +56,23 @@ class Scoreboard : public SimObject { protected: void on_reset(); + void on_tick(); private: static uint32_t get_reg_id(const RegOpd& reg, uint32_t wid) { return (wid << RegOpd::ID_BITS) | reg.id(); } + static constexpr uint32_t kReleaseDelay = 2; + + struct pending_release_t { + uint32_t wid; + RegOpd reg; + uint64_t due; + }; + std::vector> in_use_regs_; + std::deque pending_releases_; std::unordered_map owners_; std::unordered_map commit_counts_; diff --git a/sim/simx/sfu_unit.cpp b/sim/simx/sfu_unit.cpp index c6a9e969a3..de61756629 100644 --- a/sim/simx/sfu_unit.cpp +++ b/sim/simx/sfu_unit.cpp @@ -32,6 +32,7 @@ using namespace vortex; SfuUnit::SfuUnit(const SimContext& ctx, const char* name, Core* core) : FuncUnit(ctx, name, core, 6) + , warp_ctl_out(this, VX_CFG_NUM_WARPS) #ifdef VX_CFG_EXT_DXA_ENABLE , dxa_req_out(this) #endif @@ -67,8 +68,10 @@ SfuUnit::SfuUnit(const SimContext& ctx, const char* name, Core* core) { } +// Cycles from processing to the output: the result buffer and the PE +// switch's response stage. uint32_t SfuUnit::latency_of(const instr_trace_t* /*trace*/) const { - return 4; + return 2; } #ifdef VX_CFG_EXT_RTU_ENABLE @@ -406,8 +409,9 @@ void SfuUnit::on_tick() { } bool release_warp = trace->fetch_stall; + bool warp_exit = false; if (std::get_if(&trace->op_type)) { - release_warp = wctl_unit_->process(trace); + release_warp = wctl_unit_->process(trace, &warp_exit); } else if (std::get_if(&trace->op_type)) { csr_unit_->process(trace); #ifdef VX_CFG_EXT_DXA_ENABLE @@ -426,6 +430,13 @@ void SfuUnit::on_tick() { // sync-barrier, a not-yet-last barrier arrival, a deferred wspawn, or a // warp that disabled itself (tmask=0) keeps the warp parked — it is // released by the barrier/spawn machinery rather than at this commit. + // A warp-control op that does release its warp does so as it resolves; + // a join first pops the divergence stack, which takes one more cycle. + if (auto wctl_p = std::get_if(&trace->op_type); wctl_p && trace->eop + && (release_warp || warp_exit)) { + warp_ctl_out.send(WarpCtl{trace->wid, warp_exit}, (*wctl_p == WctlType::JOIN) ? 2 : 1); + release_warp = false; + } trace->resume_warp = release_warp; input.pop(); diff --git a/sim/simx/sfu_unit.h b/sim/simx/sfu_unit.h index 2730c800f0..1eb74a31f1 100644 --- a/sim/simx/sfu_unit.h +++ b/sim/simx/sfu_unit.h @@ -56,6 +56,10 @@ class SfuUnit : public FuncUnit { public: SfuUnit(const SimContext& ctx, const char* name, Core*); + // Resolved warp-control op of a stalled warp, registered toward the + // scheduler. + SimChannel warp_ctl_out; + CsrUnit& csr_unit() { return *csr_unit_; } #ifdef VX_CFG_EXT_DXA_ENABLE diff --git a/sim/simx/types.h b/sim/simx/types.h index 1f892713b5..6423d1a61a 100644 --- a/sim/simx/types.h +++ b/sim/simx/types.h @@ -1967,6 +1967,13 @@ struct GbarResume { uint32_t bar_id; }; +// A resolved warp-control op, registered from the SFU to the scheduler: the +// warp is released, or deactivated when it disabled all its threads. +struct WarpCtl { + uint32_t wid; + bool exit; +}; + // Fragment-work-distributor control messages, carried on raster-core <-> core // event links. struct FwdArm { diff --git a/sim/simx/wctl_unit.cpp b/sim/simx/wctl_unit.cpp index 9a0e5b1e22..4fbe8a531e 100644 --- a/sim/simx/wctl_unit.cpp +++ b/sim/simx/wctl_unit.cpp @@ -20,8 +20,9 @@ using namespace vortex; -bool WctlUnit::process(instr_trace_t* trace) { +bool WctlUnit::process(instr_trace_t* trace, bool* exit) { bool release_warp = trace->fetch_stall; + *exit = false; auto wctl_type = std::get(trace->op_type); auto& sched = core_->scheduler(); auto& warp = sched.warp(trace->wid); @@ -47,7 +48,12 @@ bool WctlUnit::process(instr_trace_t* trace) { next_tmask.set(t, rs1_data.at(thread_last).u & (1 << t)); } if (trace->eop) { - release_warp = core_->setTmask(trace->wid, next_tmask); + if (next_tmask.none()) { + *exit = true; + release_warp = false; + } else { + release_warp = core_->setTmask(trace->wid, next_tmask); + } } } break; case WctlType::WSPAWN: { diff --git a/sim/simx/wctl_unit.h b/sim/simx/wctl_unit.h index a542fd2c7a..799c34647e 100644 --- a/sim/simx/wctl_unit.h +++ b/sim/simx/wctl_unit.h @@ -28,9 +28,11 @@ class WctlUnit { explicit WctlUnit(Core* core) : core_(core) {} // Execute the WctlType side effects for `trace`. Returns whether the - // warp should be released after this trace's eop fires. Caller is - // responsible for the input/output channel push and the latency. - bool process(instr_trace_t* trace); + // warp should be released after this trace's eop fires; `exit` is set + // when a TMC disables every thread, leaving the warp's deactivation to + // the caller. Caller is responsible for the input/output channel push and + // the latency. + bool process(instr_trace_t* trace, bool* exit); private: Core* core_; From f1bb2f6a15805a10ce812727abcf3a5a61f21bd8 Mon Sep 17 00:00:00 2001 From: tinebp Date: Sat, 3 Oct 2026 03:05:08 -0700 Subject: [PATCH 7/8] simx: model the issue lock, per-slot dispatch queues and port-0 flush register as the RTL builds them Brings every model_parity case within the 5% default. The three cases over it: sgemv 6.38% -> 0.83%, wgmma-fedp2k-rs 6.21% -> 0.92%, wgmma-dxa-mcast -7.31% -> 3.32%. Each change models an RTL structure confirmed in the RTL source and against debug traces from both models. Issue and dispatch: - Issue lock: a warp inside a locked uop sequence (WGMMA, RTU TRACE, OM export) blocks every other warp of its issue slot, whatever the unit. The lock reopens when the last uop issues. SimX used to lock only the target unit and released the lock later, when the unit accepted the last uop. - Dispatch queues: one per issue slot, with the lowest slot served first. A slot's issue credit returns when its op leaves the queue (new ReleaseOut channel). Dispatchers bind straight to the unit inputs. SimX used to share one queue per unit and return the credit when the unit accepted the op. The queue order sets how CTAs split across warps in wgmma-fedp2k-rs. - The near-full flag reaches issue through two registers, and a refused operand-to-dispatch hand-off counts as that unit's stall. - The uop sequencer offers a macro's first uop one cycle after the macro reaches the buffer front, timed from each warp's issue history rather than instruction uuids. Release builds zero the uuids, so a uuid-based check timed differently in release and debug. TCU and pack-load macros no longer stall fetch. Units and memory: - Reduced-width ALU/FPU: one extra input register and one extra lane-gather output register. - Multiply/divide latency is 3; one result per cycle is kept. - TCU admission holds while a completing result finds an earlier one still waiting for commit. - The first data-cache port also carries cache-flush injection, so its requests take one more register. Same-bank pairs are therefore served port 1 first, which sets the eviction order under the cache thrash in sgemv. Commit and scoreboard: - The fflags/frm CSR fields are scoreboarded like registers. - A result crosses one more register before the commit arbiter. - Scoreboard release delay goes from 2 cycles to 1. FPU latency coupling (VX_config.toml): - The DPI FPU is a fast simulation stand-in for STD, so it now takes STD's FDIV/FSQRT latencies from the same toml helpers (17/17, or 32/32 with D) instead of its own 15/10. Changing STD's latency therefore changes DPI's too. DSP and FPNEW keep their own values. - SimX (STD) and rtlsim (DPI) now resolve to the same latencies. Perf baselines, regenerated with `pytest -m perf_gate --update-baselines`: - The latency change moves only kernels that divide or take square roots, and by little: raycast +0.17%, stencil3d +0.18%, draw3d +0.09%, softmax -0.01%, rt_raycast -0.14%. - The other moves were already present before this change: upstream RTL commits merged after the last recording, and catalog configs changed without a re-record. wgmma-dxa (-13.9%) and wgmma-dxa-mcast (+3.8%) were outside the 2% tolerance, so perf_gate was already failing on them. Verified on a 32-bit build: - model_parity: 31/31 within 5%, worst fp16 -4.51%. The fp16 margin is only partly explained. - perf_gate: 49/49 against the new baselines. - All 23 SimX CI test categories pass: 261/261. - Release and debug SimX builds give identical cycle counts on wgmma and draw3d. Co-Authored-By: Claude Opus 5.5 (1M context) --- VX_config.toml | 10 +- ci/baselines/perf/core.json | 56 ++++---- ci/baselines/perf/dxa.json | 40 +++--- ci/baselines/perf/graphics.json | 26 ++-- ci/baselines/perf/raytracing.json | 12 +- ci/baselines/perf/tensor.json | 18 +-- ci/baselines/perf/tensor_mx.json | 4 +- ci/baselines/perf/tensor_sp.json | 4 +- ci/baselines/perf/tensor_wg.json | 16 +-- sim/simx/alu_unit.cpp | 16 ++- sim/simx/alu_unit.h | 7 + sim/simx/core.cpp | 211 +++++++++++++++++++----------- sim/simx/decode.cpp | 27 +++- sim/simx/dispatcher.cpp | 21 +-- sim/simx/dispatcher.h | 6 +- sim/simx/instr.h | 17 +++ sim/simx/instr_trace.h | 7 +- sim/simx/mem/lsu_mem_adapter.cpp | 3 +- sim/simx/mem/lsu_mem_adapter.h | 4 + sim/simx/mem/mem_coalescer.cpp | 3 +- sim/simx/mem/mem_coalescer.h | 4 + sim/simx/om/om_unit.cpp | 3 + sim/simx/rtu/rtu_unit.cpp | 4 + sim/simx/scoreboard.cpp | 33 ++++- sim/simx/scoreboard.h | 10 +- sim/simx/sequencer.h | 3 + sim/simx/tcu/tcu_unit.cpp | 17 +++ 27 files changed, 384 insertions(+), 198 deletions(-) diff --git a/VX_config.toml b/VX_config.toml index 6cadf906da..71b9c59b73 100644 --- a/VX_config.toml +++ b/VX_config.toml @@ -121,8 +121,14 @@ fpu_dsp_vivado = "expr: $VX_CFG_FPU_TYPE_DSP and $VIVADO" # STD: FDIV_LATENCY must equal FSQRT_LATENCY (shared serializer in VX_fpu_std) VX_CFG_FMA_LATENCY = "expr: 16 if $fpu_dsp_vivado else (4 if $fpu_dsp_quartus else (12 if $VX_CFG_EXT_D_ENABLE else 8))" -VX_CFG_FDIV_LATENCY = "expr: 15 if $VX_CFG_FPU_TYPE_DPI else (16 if $VX_CFG_FPU_TYPE_FPNEW else ((32 if $VX_CFG_EXT_D_ENABLE else 17) if $VX_CFG_FPU_TYPE_STD else (15 if $fpu_dsp_quartus else (28 if $fpu_dsp_vivado else (32 if $VX_CFG_EXT_D_ENABLE else 17)))))" -VX_CFG_FSQRT_LATENCY= "expr: 10 if $VX_CFG_FPU_TYPE_DPI else (16 if $VX_CFG_FPU_TYPE_FPNEW else ((32 if $VX_CFG_EXT_D_ENABLE else 17) if $VX_CFG_FPU_TYPE_STD else (10 if $fpu_dsp_quartus else (28 if $fpu_dsp_vivado else (32 if $VX_CFG_EXT_D_ENABLE else 17)))))" +# STD timing; DPI is a fast STD and takes the same values. +fdiv_latency_std = "expr: 32 if $VX_CFG_EXT_D_ENABLE else 17" +fsqrt_latency_std = "expr: $fdiv_latency_std" +fdiv_latency_dsp = "expr: 15 if $QUARTUS else (28 if $VIVADO else (32 if $VX_CFG_EXT_D_ENABLE else 17))" +fsqrt_latency_dsp = "expr: 10 if $QUARTUS else (28 if $VIVADO else (32 if $VX_CFG_EXT_D_ENABLE else 17))" + +VX_CFG_FDIV_LATENCY = "expr: 16 if $VX_CFG_FPU_TYPE_FPNEW else ($fdiv_latency_dsp if $VX_CFG_FPU_TYPE_DSP else $fdiv_latency_std)" +VX_CFG_FSQRT_LATENCY= "expr: 16 if $VX_CFG_FPU_TYPE_FPNEW else ($fsqrt_latency_dsp if $VX_CFG_FPU_TYPE_DSP else $fsqrt_latency_std)" VX_CFG_FNCP_LATENCY = 2 VX_CFG_FCVT_LATENCY = 5 diff --git a/ci/baselines/perf/core.json b/ci/baselines/perf/core.json index 04d7bddb65..84da4d380c 100644 --- a/ci/baselines/perf/core.json +++ b/ci/baselines/perf/core.json @@ -1,23 +1,23 @@ { "core:raycast-nt16:rtlsim": { "32": { - "cycles": 567659, + "cycles": 568490, "instrs": 137486 }, "app": "raycast", "args": "", - "config_hash": "44ab43eb3398c534", + "config_hash": "6def30b1e0a43d1c", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:raycast-nt4:rtlsim": { "32": { - "cycles": 1492954, + "cycles": 1495440, "instrs": 259228 }, "app": "raycast", "args": "", - "config_hash": "21d66ac4dd47786c", + "config_hash": "276d76e1ff9f33cf", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -28,40 +28,40 @@ }, "app": "sgemm", "args": "-n128", - "config_hash": "633922118d45dc38", + "config_hash": "ee0243cf40ad66eb", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=16 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, "core:sgemm-mc-nt4:rtlsim": { "32": { - "cycles": 2206450, + "cycles": 2211340, "instrs": 2494528 }, "app": "sgemm", "args": "-n128", - "config_hash": "f8c8f3e3090e2f7b", + "config_hash": "ff3f840c868f11fc", "configs": "-DVX_CFG_NUM_WARPS=8 -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, "core:sgemm-nt16:rtlsim": { "32": { - "cycles": 1821882, + "cycles": 1823521, "instrs": 623632 }, "app": "sgemm", "args": "-n128", - "config_hash": "73ae57f76a7963ec", + "config_hash": "54d7660022bc0244", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:sgemm-nt4:rtlsim": { "32": { - "cycles": 5957013, + "cycles": 6055079, "instrs": 2494480 }, "app": "sgemm", "args": "-n128", - "config_hash": "c9f6f1d6a54550ef", + "config_hash": "93c5b85b94bf8538", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -72,18 +72,18 @@ }, "app": "sgemmx", "args": "-n128", - "config_hash": "2eb8c2e7ab22c3cd", + "config_hash": "29d1dfbed14e9155", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:sgemmx-nt4:rtlsim": { "32": { - "cycles": 4292070, + "cycles": 4210203, "instrs": 1067536 }, "app": "sgemmx", "args": "-n128", - "config_hash": "ec8b33c4a7153f3f", + "config_hash": "9f17cd000d55867e", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -94,7 +94,7 @@ }, "app": "sgemv", "args": "-m512 -n512", - "config_hash": "50c66d5a22a7e8ba", + "config_hash": "517d45e208f26458", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -105,73 +105,73 @@ }, "app": "sgemv", "args": "-m512 -n512", - "config_hash": "60546fb16d4af21c", + "config_hash": "786b7e475b0e7ee9", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, "core:softmax-nt16:rtlsim": { "32": { - "cycles": 1274529, + "cycles": 1274656, "instrs": 464089 }, "app": "softmax", "args": "-n64", - "config_hash": "9404c6ac13e2b572", + "config_hash": "46be2b39ddbf6aa5", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:softmax-nt4:rtlsim": { "32": { - "cycles": 4889616, + "cycles": 4860631, "instrs": 1789797 }, "app": "softmax", "args": "-n64", - "config_hash": "ef5d2c0ed37cab34", + "config_hash": "2d2af3b50047fd34", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, "core:stencil3d-nt16:rtlsim": { "32": { - "cycles": 2384949, + "cycles": 2388738, "instrs": 892944 }, "app": "stencil3d", "args": "-n32", - "config_hash": "18b8564a0ad8446b", + "config_hash": "83f246361b1a2b62", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:stencil3d-nt4:rtlsim": { "32": { - "cycles": 4195670, + "cycles": 4177542, "instrs": 1785872 }, "app": "stencil3d", "args": "-n32", - "config_hash": "fec853f7b4a7fa5e", + "config_hash": "76870db61c0f4de7", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, "core:vecadd-nt16:rtlsim": { "32": { - "cycles": 75335, + "cycles": 75714, "instrs": 24592 }, "app": "vecadd", "args": "-n16384", - "config_hash": "c218290fa961b753", + "config_hash": "9b4c96e327b73aba", "configs": "-DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "core:vecadd-nt4:rtlsim": { "32": { - "cycles": 292046, + "cycles": 293078, "instrs": 98320 }, "app": "vecadd", "args": "-n16384", - "config_hash": "0b04aaed6e773896", + "config_hash": "ad2cd42528d559e6", "configs": "-DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" } diff --git a/ci/baselines/perf/dxa.json b/ci/baselines/perf/dxa.json index c83b5b070d..14af4af964 100644 --- a/ci/baselines/perf/dxa.json +++ b/ci/baselines/perf/dxa.json @@ -6,7 +6,7 @@ }, "app": "dxa_copy_mcast", "args": "", - "config_hash": "0aff16ac6b8499b4", + "config_hash": "039e26d60791f884", "configs": "-DVX_CFG_EXT_DXA_ENABLE", "driver": "rtlsim" }, @@ -17,7 +17,7 @@ }, "app": "dxa_copy", "args": "-d1", - "config_hash": "f0e7503341d9eac3", + "config_hash": "ab44b87188699391", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -28,29 +28,29 @@ }, "app": "dxa_copy", "args": "-d1", - "config_hash": "34210b7d0de2ad82", + "config_hash": "40c3e5367ee82910", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, "dxa:perf_gate-sgemm2-dxa-mcast-nt16:rtlsim": { "32": { - "cycles": 14332355, + "cycles": 14332387, "instrs": 4961296 }, "app": "sgemm2_dxa_mcast", "args": "-n128", - "config_hash": "1b5a3bdc898873dc", + "config_hash": "7556f3195acb39b1", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "dxa:perf_gate-sgemm2-dxa-mcast-nt4:rtlsim": { "32": { - "cycles": 14344871, + "cycles": 14344887, "instrs": 4961296 }, "app": "sgemm2_dxa_mcast", "args": "-n128", - "config_hash": "5a08ea62af7bf678", + "config_hash": "684af59553c0dd7b", "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": "3a732a3d8f3aaf1f", + "config_hash": "1b903f48fc402adf", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, @@ -72,41 +72,41 @@ }, "app": "sgemm2_dxa", "args": "-n128 -t4 -m1", - "config_hash": "edd308f5092a3c9a", + "config_hash": "dbdacf51150051f7", "configs": "-DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, "dxa:perf_gate-wgmma-dxa-mcast-nt16:rtlsim": { "32": { - "cycles": 108106, + "cycles": 99299, "instrs": 32384 }, "app": "sgemm_tcu_wg_dxa_mcast", "args": "-m 128 -n 128 -k 64", - "config_hash": "94ff36ea84f276ed", - "configs": "-DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=16", + "config_hash": "d0cb73c9d6f7ffaa", + "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" }, "dxa:perf_gate-wgmma-dxa-mcast-nt4:rtlsim": { "32": { - "cycles": 393379, + "cycles": 408190, "instrs": 234560 }, "app": "sgemm_tcu_wg_dxa_mcast", "args": "-m 128 -n 128 -k 64", - "config_hash": "5f35238719d0a258", - "configs": "-DVX_CFG_EXT_TCU_ENABLE -DVX_CFG_EXT_DXA_ENABLE -DVX_CFG_NUM_THREADS=4", + "config_hash": "e767bf6bcc2b1d4a", + "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" }, "dxa:perf_gate-wgmma-dxa:rtlsim": { "32": { - "cycles": 46211, + "cycles": 39805, "instrs": 9160 }, "app": "sgemm_tcu_wg_dxa", "args": "-m 128 -n 128 -k 64", - "config_hash": "84fdf91752259a37", - "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_WARPS=8 -DVX_CFG_NUM_THREADS=32", + "config_hash": "271a7cf569c588e4", + "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" }, "dxa:perf_gate-wgmma-sp-dxa:rtlsim": { @@ -116,8 +116,8 @@ }, "app": "sgemm_tcu_wg_sp_dxa", "args": "-m 128 -n 128 -k 128", - "config_hash": "528ad9e91be992e1", - "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_WARPS=8 -DVX_CFG_NUM_THREADS=32", + "config_hash": "1db9ac61c1e7ec5d", + "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 6e35026942..39103db55e 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": 531412, + "cycles": 531658, "instrs": 204102 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "3decf2ea51277f28", + "config_hash": "0faab87874bca729", "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": 149499, + "cycles": 148046, "instrs": 51475 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "fe4f8017dee4867a", + "config_hash": "a553842e84b4fa29", "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": 532663, + "cycles": 531210, "instrs": 204102 }, "app": "gfx_draw3d", "args": "-w 128 -h 128 -t box.cgltrace -r box_ref_128.png", - "config_hash": "711212427ad45bc3", + "config_hash": "e28dc991c4fcd5df", "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": "ed94f2a729c96a45", + "config_hash": "1dc6dd362b5bc045", "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": "df22122a0e696f68", + "config_hash": "8b7ac7fc4d0d0d78", "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": "eefa3aeb3bc1b0d6", + "config_hash": "22b248ec7b5ce453", "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": "be8529070b7e847a", + "config_hash": "f99789ca113535d5", "configs": "-DVX_CFG_EXT_RASTER_ENABLE -DVX_CFG_NUM_THREADS=4", "driver": "rtlsim" }, @@ -83,18 +83,18 @@ }, "app": "gfx_tex", "args": "-i toad.png -r toad_ref_f0.png -f 0 -g 0", - "config_hash": "452d34e0fe700df1", + "config_hash": "2260e1daed5f2e6c", "configs": "-DVX_CFG_EXT_TEX_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "graphics:perf_gate-tex-nt4:rtlsim": { "32": { - "cycles": 137275, + "cycles": 136738, "instrs": 49168 }, "app": "gfx_tex", "args": "-i toad.png -r toad_ref_f0.png -f 0 -g 0", - "config_hash": "b0c8d09eeb0cdbf7", + "config_hash": "29b6aaeb05512055", "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 3c1c88430a..46d0cbece7 100644 --- a/ci/baselines/perf/raytracing.json +++ b/ci/baselines/perf/raytracing.json @@ -1,34 +1,34 @@ { "raytracing:perf_gate-rt_raycast-mc:rtlsim": { "32": { - "cycles": 82712, + "cycles": 85157, "instrs": 39216 }, "app": "rt_raycast", "args": "", - "config_hash": "b0bdc8a6ee3db951", + "config_hash": "f1423c026ac8b273", "configs": "-DVX_CFG_EXT_RTU_ENABLE -DVX_CFG_NUM_CORES=2 -DVX_CFG_NUM_THREADS=4 -DVX_CFG_L2_ENABLE", "driver": "rtlsim" }, "raytracing:perf_gate-rt_raycast-nt16:rtlsim": { "32": { - "cycles": 51652, + "cycles": 52605, "instrs": 11310 }, "app": "rt_raycast", "args": "", - "config_hash": "3417d94d526bfe77", + "config_hash": "e8694976d2b3d2b5", "configs": "-DVX_CFG_EXT_RTU_ENABLE -DVX_CFG_NUM_THREADS=16", "driver": "rtlsim" }, "raytracing:perf_gate-rt_raycast-nt4:rtlsim": { "32": { - "cycles": 156410, + "cycles": 158049, "instrs": 39200 }, "app": "rt_raycast", "args": "", - "config_hash": "5e5e62aa53e4c487", + "config_hash": "52ba7591a71b9649", "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 fb9764766a..30c6a530d3 100644 --- a/ci/baselines/perf/tensor.json +++ b/ci/baselines/perf/tensor.json @@ -6,8 +6,8 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "af0310169ce2509a", - "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", + "config_hash": "0a61c4b9cadb1d3e", + "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" }, "tensor:perf_gate-fp16-mc-nt16:rtlsim": { @@ -17,18 +17,18 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "f3eda425ced1475d", + "config_hash": "8691a66768ce4d5a", "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" }, "tensor:perf_gate-fp16-mc-nt4:rtlsim": { "32": { - "cycles": 222045, + "cycles": 221331, "instrs": 336448 }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "2d05619451c35ff6", + "config_hash": "30dd902fdb9a560c", "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,19 +39,19 @@ }, "app": "sgemm_tcu", "args": "-m 128 -n 128 -k 128", - "config_hash": "49026bd95d5a1817", + "config_hash": "91b77b52210fc6cd", "configs": "-DVX_CFG_NUM_THREADS=8 -DVX_CFG_EXT_TCU_ENABLE -DITYPE=fp16 -DOTYPE=fp32", "driver": "rtlsim" }, "tensor:perf_gate-sgemm2-fp16:rtlsim": { "32": { - "cycles": 601768, + "cycles": 592658, "instrs": 409888 }, "app": "sgemm2_tcu", "args": "-m 128 -n 128 -k 64", - "config_hash": "0f32f92417fbaa3c", - "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", + "config_hash": "de3d4ecc5d7bdce7", + "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 3695e06ef7..c1475de5ab 100644 --- a/ci/baselines/perf/tensor_mx.json +++ b/ci/baselines/perf/tensor_mx.json @@ -1,12 +1,12 @@ { "tensor_mx:perf_gate-mxfp8:rtlsim": { "32": { - "cycles": 77942, + "cycles": 77127, "instrs": 34832 }, "app": "sgemm_tcu_mx", "args": "-m 128 -n 128 -k32", - "config_hash": "e6d3cb245141f6af", + "config_hash": "74c641d0b30e9646", "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 e4a3027be3..16e50b8a12 100644 --- a/ci/baselines/perf/tensor_sp.json +++ b/ci/baselines/perf/tensor_sp.json @@ -1,12 +1,12 @@ { "tensor_sp:perf_gate-sparse-fp16:rtlsim": { "32": { - "cycles": 188923, + "cycles": 187290, "instrs": 100880 }, "app": "sgemm_tcu_sp", "args": "-m 128 -n 128 -k 128", - "config_hash": "1426278a05ab7a51", + "config_hash": "e270a1261dd01a09", "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 2e77cfa966..854e92aa56 100644 --- a/ci/baselines/perf/tensor_wg.json +++ b/ci/baselines/perf/tensor_wg.json @@ -1,13 +1,13 @@ { "tensor_wg:perf_gate-wgmma-fedp2k-rs:rtlsim": { "32": { - "cycles": 159683, + "cycles": 157272, "instrs": 108512 }, "app": "sgemm_tcu_wg", "args": "-m64 -n64 -k64", - "config_hash": "84f3ab0ee89d1c7e", - "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_FEDP2K -DITYPE=fp16 -DOTYPE=fp32 -DWGMMA_NRC=8 -DWGMMA_RS", + "config_hash": "5098d4c8b863f240", + "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" }, "tensor_wg:perf_gate-wgmma-fp16-ss:rtlsim": { @@ -17,19 +17,19 @@ }, "app": "sgemm_tcu_wg", "args": "-m 128 -n 128 -k 128", - "config_hash": "51c5e22e5cab34a8", - "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", + "config_hash": "2b1479f2108a57cc", + "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" }, "tensor_wg:perf_gate-wgmma-sparse:rtlsim": { "32": { - "cycles": 408016, + "cycles": 401322, "instrs": 258976 }, "app": "sgemm_tcu_wg_sp", "args": "-m 128 -n 128 -k 128", - "config_hash": "b94d94645f0461c8", - "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", + "config_hash": "f96b004699e182e0", + "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/sim/simx/alu_unit.cpp b/sim/simx/alu_unit.cpp index 4da3a189fd..a866e2f1ca 100644 --- a/sim/simx/alu_unit.cpp +++ b/sim/simx/alu_unit.cpp @@ -25,7 +25,10 @@ using namespace vortex; AluUnit::AluUnit(const SimContext& ctx, const char* name, Core* core) - : FuncUnit(ctx, name, core) + // The output channel also counts results still in flight: it covers the + // deepest path plus the result waiting for the commit side, so the + // multiplier keeps one result per cycle. + : FuncUnit(ctx, name, core, kMulDivLatency + (kGather ? 1 : 0) + 2) , branch_ctl_out(this, VX_CFG_NUM_WARPS) {} @@ -73,9 +76,9 @@ uint32_t AluUnit::latency_of(const instr_trace_t* trace) const { } else if (std::get_if(&trace->op_type)) { auto mdv_type = std::get(trace->op_type); switch (mdv_type) { - // The multiplier pipeline is three stages deep against the integer - // ALU's single response stage; simulation divides run in that same - // pipeline rather than iteratively. + // Three multiplier stages plus the multiply/divide response register, + // against the integer ALU's single response stage; simulation divides + // run in that same pipeline rather than iteratively. case MdvType::MUL: case MdvType::MULHU: case MdvType::MULH: @@ -84,7 +87,7 @@ uint32_t AluUnit::latency_of(const instr_trace_t* trace) const { case MdvType::DIVU: case MdvType::REM: case MdvType::REMU: - return 2; + return kMulDivLatency; default: std::abort(); } @@ -572,8 +575,7 @@ void AluUnit::on_tick() { branch_ctl_out.send(trace->wid, 1); trace->resume_warp = false; } - uint32_t delay = this->latency_of(trace); - output.send(trace, delay); + output.send(trace, this->latency_of(trace) + (kGather ? 1 : 0)); input.pop(); } } diff --git a/sim/simx/alu_unit.h b/sim/simx/alu_unit.h index b94904cc97..6e2db38121 100644 --- a/sim/simx/alu_unit.h +++ b/sim/simx/alu_unit.h @@ -19,6 +19,13 @@ namespace vortex { class AluUnit : public FuncUnit { public: + // Cycles the multiply/divide pipeline adds over the integer result stage. + static constexpr uint32_t kMulDivLatency = 3; + + // Partial bandwidth: a lane-gather stage registers the results once more. + static constexpr bool kGather = (VX_CFG_NUM_ALU_BLOCKS != VX_CFG_ISSUE_WIDTH) + || (VX_CFG_NUM_ALU_LANES != VX_CFG_SIMD_WIDTH); + AluUnit(const SimContext& ctx, const char* name, Core*); // Resolved branch of a stalled warp, registered toward the scheduler. diff --git a/sim/simx/core.cpp b/sim/simx/core.cpp index cdaf880714..f67dba6e23 100644 --- a/sim/simx/core.cpp +++ b/sim/simx/core.cpp @@ -70,10 +70,13 @@ class Core::Impl { , decode_latch_(ctx, "decode_latch", 1, 2) , pending_icache_(VX_CFG_NUM_WARPS) , ibuffer_arbs_(VX_CFG_ISSUE_WIDTH, {ArbiterType::GTO, PER_ISSUE_WARPS}) - , fu_locked_(VX_CFG_ISSUE_WIDTH, BitVector<>((uint32_t)FUType::Count, 0)) - , fu_unlock_pending_(VX_CFG_ISSUE_WIDTH, std::vector((uint32_t)FUType::Count, nullptr)) + , issue_lock_owner_(VX_CFG_ISSUE_WIDTH, kLockOpen) , fu_credits_(VX_CFG_ISSUE_WIDTH, std::vector((uint32_t)FUType::Count, 0)) + , fu_full_(VX_CFG_ISSUE_WIDTH, BitVector<>((uint32_t)FUType::Count, 0)) + , fu_full_seen_(VX_CFG_ISSUE_WIDTH, BitVector<>((uint32_t)FUType::Count, 0)) , ibuf_inflight_(VX_CFG_NUM_WARPS, 0) + , last_issue_(VX_CFG_NUM_WARPS, 0) + , staged_since_(VX_CFG_NUM_WARPS, 0) { const std::string& name = simobject_->name(); char sname[100]; @@ -165,6 +168,8 @@ class Core::Impl { std::vector*> dc_req_out(VX_CFG_NUM_LSU_BLOCKS * DCACHE_CHANNELS); std::vector*> dc_rsp_in(VX_CFG_NUM_LSU_BLOCKS * DCACHE_CHANNELS); + // The first data-cache port also carries the cache-flush injection, which + // registers its requests once more. if ((VX_CFG_NUM_LSU_LANES > 1) && (DCACHE_WORD_SIZE > LSU_WORD_SIZE)) { // connect memory coalescer; its memory side drives the dcache // channels directly (combinational lane fan-out/fan-in). @@ -176,6 +181,7 @@ class Core::Impl { dc_rsp_in.at(b * DCACHE_CHANNELS + c) = &mem_coalescers_.at(b)->RspIn.at(c); } } + mem_coalescers_.at(0)->set_port_delay(0, 1); } else { // bypass memory coalescer: per-block lane adapter (channel-fused // pass-through when DCACHE_CHANNELS == 1) @@ -190,6 +196,7 @@ class Core::Impl { dc_rsp_in.at(b * DCACHE_CHANNELS + c) = &lsu_dcache_adapter.at(b)->RspIn.at(c); } } + lsu_dcache_adapter.at(0)->set_port_delay(0, 1); } #ifdef VX_CFG_VM_ENABLE @@ -229,15 +236,6 @@ class Core::Impl { } #endif - // initialize dispatchers - dispatchers_.at((int)FUType::ALU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_ALU_BLOCKS, VX_CFG_NUM_ALU_LANES); - dispatchers_.at((int)FUType::FPU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_FPU_BLOCKS, VX_CFG_NUM_FPU_LANES); - dispatchers_.at((int)FUType::LSU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_LSU_BLOCKS, VX_CFG_NUM_LSU_LANES); - dispatchers_.at((int)FUType::SFU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_SFU_BLOCKS, VX_CFG_NUM_SFU_LANES); - #ifdef VX_CFG_EXT_TCU_ENABLE - dispatchers_.at((int)FUType::TCU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_TCU_BLOCKS, VX_CFG_NUM_TCU_LANES); - #endif - // initialize execute units snprintf(sname, 100, "%s-alu", name.c_str()); func_units_.at((int)FUType::ALU) = SimPlatform::instance().create_object(sname, simobject_); @@ -276,6 +274,28 @@ class Core::Impl { #endif #endif + // initialize dispatchers after the units they feed, so a unit takes its + // input before the dispatcher fills it in the same cycle. + { + // A reduced-width ALU or FPU registers the op once more on its way in. + uint32_t alu_dly = AluUnit::kGather ? 2 : 1; + uint32_t fpu_dly = FpuUnit::kGather ? 2 : 1; + dispatchers_.at((int)FUType::ALU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_ALU_BLOCKS, VX_CFG_NUM_ALU_LANES, alu_dly); + dispatchers_.at((int)FUType::FPU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_FPU_BLOCKS, VX_CFG_NUM_FPU_LANES, fpu_dly); + dispatchers_.at((int)FUType::LSU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_LSU_BLOCKS, VX_CFG_NUM_LSU_LANES, 1); + dispatchers_.at((int)FUType::SFU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_SFU_BLOCKS, VX_CFG_NUM_SFU_LANES, 1); + #ifdef VX_CFG_EXT_TCU_ENABLE + dispatchers_.at((int)FUType::TCU) = SimPlatform::instance().create_object(name.c_str(), simobject_, VX_CFG_DISPATCH_QUEUE_SIZE, VX_CFG_NUM_TCU_BLOCKS, VX_CFG_NUM_TCU_LANES, 1); + #endif + for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { + auto& dispatch = dispatchers_.at(fu); + auto& func_unit = func_units_.at(fu); + for (uint32_t b = 0; b < func_unit->num_blocks(); ++b) { + dispatch->Outputs.at(b).bind(&func_unit->input(b)); + } + } + } + // commit queues — per-iw, per-FU staging fed at runtime in commit() by // routing per-block FU outputs on trace->wid (no static binding because // the iw is not knowable at setup time when NUM_*_BLOCKS < @@ -284,7 +304,9 @@ class Core::Impl { for (uint32_t iw = 0; iw < VX_CFG_ISSUE_WIDTH; ++iw) { auto& queues = commit_queues_.emplace_back(); for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { - queues.emplace_back(std::make_unique>(simobject_, 2)); + // Covers the results still crossing into the queue, so a unit can + // retire one result per cycle. + queues.emplace_back(std::make_unique>(simobject_, 3)); } } @@ -298,12 +320,13 @@ class Core::Impl { for (auto& fc : fu_credits_) { std::fill(fc.begin(), fc.end(), 0); } - for (auto& fl : fu_locked_) { - fl.reset(); - } - for (auto& fp : fu_unlock_pending_) { - std::fill(fp.begin(), fp.end(), nullptr); + for (uint32_t iw = 0; iw < VX_CFG_ISSUE_WIDTH; ++iw) { + fu_full_.at(iw).reset(); + fu_full_seen_.at(iw).reset(); } + std::fill(issue_lock_owner_.begin(), issue_lock_owner_.end(), kLockOpen); + std::fill(last_issue_.begin(), last_issue_.end(), 0); + std::fill(staged_since_.begin(), staged_since_.end(), 0); pending_instrs_.clear(); pending_ifetches_ = 0; @@ -319,8 +342,9 @@ class Core::Impl { void tick() { this->commit(); - this->execute(); + this->return_dispatch_credits(); this->issue(); + this->update_queue_full(); this->decode(); this->fetch(); this->schedule(); @@ -522,6 +546,7 @@ class Core::Impl { DT(3, simobject_->name() << "-pipeline decode: " << *trace); // insert to ibuffer + trace->ibuf_time = SimPlatform::instance().cycles(); ibuffer->push(trace); decode_latch_.pop(); @@ -534,8 +559,20 @@ class Core::Impl { if (operand->Output.empty()) continue; auto trace = operand->Output.peek(); - if (dispatchers_.at((int)trace->fu_type)->Inputs.at(iw).try_send(trace)) { + // The collector's output register adds a cycle before the slot queue. + if (dispatchers_.at((int)trace->fu_type)->Inputs.at(iw).try_send(trace, 2)) { operand->Output.pop(); + } else { + switch (trace->fu_type) { + case FUType::ALU: ++perf_stats_.alu_stalls; break; + case FUType::FPU: ++perf_stats_.fpu_stalls; break; + case FUType::LSU: ++perf_stats_.lsu_stalls; break; + case FUType::SFU: ++perf_stats_.sfu_stalls; break; + #ifdef VX_CFG_EXT_TCU_ENABLE + case FUType::TCU: ++perf_stats_.tcu_stalls; break; + #endif + default: assert(false); + } } } @@ -554,6 +591,16 @@ class Core::Impl { auto seq = sequencers_.at(wid); auto uop_trace = seq->get(trace); // returns cached uop or generates next + // The uop sequencer starts a macro-op once it reaches the front of the + // buffer behind the staged instruction, and offers its first uop the + // cycle after. A macro that arrives in an empty buffer, or behind an + // instruction that issues at once, waits that extra cycle. + bool seq_starting = false; + if (seq->starting()) { + uint64_t ready = std::max(staged_since_.at(wid), trace->ibuf_time) + 2; + seq_starting = (SimPlatform::instance().cycles() < ready); + } + if (scoreboard_->in_use(uop_trace)) { auto uses = scoreboard_->get_uses(uop_trace); if (!uop_trace->log_once(true)) { @@ -571,17 +618,20 @@ class Core::Impl { any_scrb_blocked = true; } else { uop_trace->log_once(false); - // FU lock: block warps whose target FU is locked by another warp. - // fu_lock=1 means acquire request; blocked when FU already locked. - auto fu = (int)uop_trace->fu_type; - bool uop_fu_lock = uop_trace->instr_ptr->get_fu_lock(); - if (fu_locked_.at(iw).test(fu) && uop_fu_lock) { - continue; // blocked by FU lock + if (seq_starting) { + continue; + } + // While a warp holds the issue lock, no other warp of the slot issues, + // whatever its unit. + auto lock_owner = issue_lock_owner_.at(iw); + if (lock_owner != kLockOpen && lock_owner != w) { + continue; } - // FU dispatch queue going-full: the warp does not request. Credits also - // count ops still in operand collection; the one-slot guard band keeps - // an issued op from blocking the shared operand path. - if (fu_credits_.at(iw).at(fu) >= VX_CFG_DISPATCH_QUEUE_SIZE - 1) { + auto fu = (int)uop_trace->fu_type; + // FU dispatch queue near full: the warp does not request. Credits also + // count ops still in operand collection. The flag reaches issue through + // two registers, so a warp sees the count as of two cycles back. + if (fu_full_seen_.at(iw).test(fu)) { continue; } #ifdef VX_CFG_EXT_RTU_ENABLE @@ -615,31 +665,37 @@ class Core::Impl { operands_.at(iw)->fetch_operands(uop_trace); // spend a dispatch credit for the target FU ++fu_credits_.at(iw).at((int)uop_trace->fu_type); + // An issue candidate is staged from the cycle after both the warp's + // previous issue and its own arrival in the buffer. + staged_since_.at(wid) = std::max(last_issue_.at(wid), trace->ibuf_time) + 1; + last_issue_.at(wid) = SimPlatform::instance().cycles(); DT(3, simobject_->name() << "-pipeline issue: " << *uop_trace); if (uop_trace->wb) { // update scoreboard scoreboard_->reserve(uop_trace); } - // Update FU lock state: 10=acquire, 01=release. The lock keeps a - // uop sequence contiguous at its functional unit, but the operand - // collectors and their arbiter can reorder warps between issue - // and the unit's input, so the release is deferred until the unit - // accepts the sequence's last uop (see execute()). + if (uop_trace->instr_ptr->fcsr_writes()) { + scoreboard_->reserve_fcsr(uop_trace); + } + // Issue lock: a sequence's first uop (lock without unlock) takes it + // for its warp, the last (unlock) reopens the slot from the next cycle. { - auto fui = (int)uop_trace->fu_type; bool fl = uop_trace->instr_ptr->get_fu_lock(); bool ful = uop_trace->instr_ptr->get_fu_unlock(); if (fl && !ful) { - fu_locked_.at(iw).set(fui); - } else if (!fl && ful) { - fu_unlock_pending_.at(iw).at(fui) = uop_trace; + issue_lock_owner_.at(iw) = w; + } else if (ful) { + issue_lock_owner_.at(iw) = kLockOpen; } } // Advance sequencer; pop ibuffer only when all micro-ops issued if (seq->advance()) { - // Resume warp for macro instructions that stalled fetch at decode if (trace->instr_ptr->is_macro_op()) { - scheduler_->resume(trace->wid); + // A macro that stalled fetch at decode releases its warp once + // its last micro-op issues. + if (trace->instr_ptr->is_wstall()) { + scheduler_->resume(trace->wid); + } // Macro trace never reaches commit (only micro-ops do), // so remove it from pending tracking and deallocate here. pending_instrs_.remove(trace); @@ -671,47 +727,32 @@ class Core::Impl { } } - void execute() { - // Dispatcher.Outputs are sized per FU's NUM_*_BLOCKS; FU.Inputs match. - // Per-block 1:1 forward (the dispatcher already handled IW→NB aggregation). - for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { - auto& dispatch = dispatchers_.at(fu); - auto& func_unit = func_units_.at(fu); - uint32_t nb = func_unit->num_blocks(); - for (uint32_t b = 0; b < nb; ++b) { - if (dispatch->Outputs.at(b).empty()) - continue; - auto trace = dispatch->Outputs.at(b).peek(); - if (func_unit->input(b).try_send(trace)) { - dispatch->Outputs.at(b).pop(); - // return the dispatch credit on FU accept - uint32_t iw = trace->wid % VX_CFG_ISSUE_WIDTH; - if (fu_credits_.at(iw).at(fu) > 0) - --fu_credits_.at(iw).at(fu); - // The sequence's last uop reached the unit: release its FU lock. - // A pid-split sequence hands the original trace over last. - auto& pending_unlock = fu_unlock_pending_.at(iw).at(fu); - if (pending_unlock == trace) { - fu_locked_.at(iw).reset(fu); - pending_unlock = nullptr; - } + void update_queue_full() { + for (uint32_t iw = 0; iw < VX_CFG_ISSUE_WIDTH; ++iw) { + fu_full_seen_.at(iw) = fu_full_.at(iw); + for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { + bool full = (fu_credits_.at(iw).at(fu) >= VX_CFG_DISPATCH_QUEUE_SIZE - 1); + if (full) { + fu_full_.at(iw).set(fu); } else { - // track functional unit stalls - switch ((FUType)fu) { - case FUType::ALU: ++perf_stats_.alu_stalls; break; - case FUType::FPU: ++perf_stats_.fpu_stalls; break; - case FUType::LSU: ++perf_stats_.lsu_stalls; break; - case FUType::SFU: ++perf_stats_.sfu_stalls; break; - #ifdef VX_CFG_EXT_TCU_ENABLE - case FUType::TCU: ++perf_stats_.tcu_stalls; break; - #endif - default: assert(false); - } + fu_full_.at(iw).reset(fu); } } } } + void return_dispatch_credits() { + for (uint32_t fu = 0; fu < (uint32_t)FUType::Count; ++fu) { + auto& release = dispatchers_.at(fu)->ReleaseOut; + while (!release.empty()) { + uint32_t iw = release.peek()->wid % VX_CFG_ISSUE_WIDTH; + assert(fu_credits_.at(iw).at(fu) > 0); + --fu_credits_.at(iw).at(fu); + release.pop(); + } + } + } + void commit() { // A resolved branch or warp-control op releases its warp, or retires it // when it disabled all its threads, a cycle after it resolves; the @@ -748,7 +789,8 @@ class Core::Impl { auto trace = fu_out.peek(); uint32_t iw = trace->wid % VX_CFG_ISSUE_WIDTH; auto& arb_in = *commit_queues_.at(iw).at(fu); - if (arb_in.try_send(trace)) { + // A unit registers its result once more before the commit arbiter. + if (arb_in.try_send(trace, 2)) { // Release the warp as soon as its stalling instruction's result leaves // the functional unit — the branch target / fence / warp-control is // resolved at that point. The release lands in stalled_warps so the @@ -812,6 +854,9 @@ class Core::Impl { scoreboard_->release(trace); } } + if (trace->eop && trace->instr_ptr->fcsr_writes()) { + scoreboard_->release_fcsr(trace); + } if (trace->eop) { @@ -886,9 +931,13 @@ class Core::Impl { ibuffer->pop(); } ibuf_inflight_.at(wid) = 0; + uint32_t iw = wid % VX_CFG_ISSUE_WIDTH; + if (issue_lock_owner_.at(iw) == wid / VX_CFG_ISSUE_WIDTH) { + issue_lock_owner_.at(iw) = kLockOpen; + } // The sequencer may cache the just-flushed trace in state_.current_uop // (set by seq->get() during a prior issue tick where the trace stalled - // on scoreboard or FU lock). That cached pointer is now dangling — + // on scoreboard or the issue lock). That cached pointer is now dangling — // drop it so the post-trap issue cycle re-derives state from the // post-mret ibuffer. sequencers_.at(wid)->flush(); @@ -1059,11 +1108,15 @@ class Core::Impl { std::vector ibuffer_arbs_; - std::vector> fu_locked_; - std::vector> fu_unlock_pending_; // [iw][fu] last uop of a locked sequence, released on FU accept + static constexpr uint32_t kLockOpen = ~0u; + std::vector issue_lock_owner_; // [iw] slot-local warp holding the issue lock std::vector> fu_credits_; // [iw][fu] in-flight dispatch credits + std::vector> fu_full_; // [iw] credits at or past the near-full mark + std::vector> fu_full_seen_; // [iw] that flag one cycle later, as issue sees it std::vector ibuf_inflight_; + std::vector last_issue_; // [wid] cycle of the warp's last issue + std::vector staged_since_; // [wid] cycle the last issued candidate was staged PoolAllocator trace_pool_; diff --git a/sim/simx/decode.cpp b/sim/simx/decode.cpp index 6d4fcc2120..9fb835db56 100644 --- a/sim/simx/decode.cpp +++ b/sim/simx/decode.cpp @@ -677,6 +677,17 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { std::abort(); } auto imm12 = code >> shift_rs2; + { + bool csr_write = (funct3 == 1) || (funct3 == 5) || (rs1 != 0); + uint8_t fcsr_regs = 0; + if (imm12 == VX_CSR_FFLAGS || imm12 == VX_CSR_FCSR) { + fcsr_regs |= Instr::FCSR_FFLAGS; + } + if (imm12 == VX_CSR_FRM || imm12 == VX_CSR_FCSR) { + fcsr_regs |= Instr::FCSR_FRM; + } + instr->set_fcsr_use(fcsr_regs, csr_write ? fcsr_regs : 0); + } if (funct3 < 5) { instr->set_src_reg(0, rs1, RegType::Integer); instr->set_args(IntrCsrArgs{0, 0, imm12}); @@ -693,6 +704,8 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { case Opcode::FCI: { instr->set_fu_type(FUType::FPU); instr->set_args(IntrFpuArgs{funct3, rs2, (funct7 & 0x1)}); + // Most FP ops raise exception flags; a dynamic rounding mode reads frm. + uint8_t frm_rd = (funct3 == 0x7) ? Instr::FCSR_FRM : 0; switch (funct7) { case 0x00: // RV32F: FADD.S case 0x01: // RV32D: FADD.D @@ -700,6 +713,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x04: // RV32F: FSUB.S case 0x05: // RV32D: FSUB.D @@ -707,6 +721,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x08: // RV32F: FMUL.S case 0x09: // RV32D: FMUL.D @@ -714,6 +729,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x10: // RV32F: FSGNJ.S, FSGNJN.S, FSGNJX.S case 0x11: // RV32D: FSGNJ.D, FSGNJN.D, FSGNJX.D @@ -728,6 +744,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(0, Instr::FCSR_FFLAGS); break; case 0x0c: // RV32F: FDIV.S case 0x0d: // RV32D: FDIV.D @@ -735,18 +752,21 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x20: // FCVT.S.D case 0x21: // FCVT.D.S instr->set_op_type(FpuType::F2F); instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x2c: // FSQRT.S case 0x2d: // FSQRT.D instr->set_op_type(FpuType::FSQRT); instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Float); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x50: // FLE.S, FLT.S, FEQ.S case 0x51: // FLE.D, FLT.D, FEQ.D @@ -754,6 +774,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Integer); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); + instr->set_fcsr_use(0, Instr::FCSR_FFLAGS); break; case 0x60: // FCVT.W.D, FCVT.WU.D, FCVT.L.D, FCVT.LU.D case 0x61: // FCVT.W.S, FCVT.WU.S, FCVT.L.S, FCVT.LU.S @@ -761,6 +782,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Integer); instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::None); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x68: // FCVT.S.W, FCVT.S.WU, FCVT.S.L, FCVT.S.LU case 0x69: // FCVT.D.W, FCVT.D.WU, FCVT.D.L, FCVT.D.LU @@ -768,6 +790,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_dest_reg(rd, RegType::Float); instr->set_src_reg(0, rs1, RegType::Integer); instr->set_src_reg(1, rs2, RegType::None); + instr->set_fcsr_use(frm_rd, Instr::FCSR_FFLAGS); break; case 0x70: // FCLASS.S, FMV.X.S case 0x71: // FCLASS.D, FMV.X.D @@ -798,6 +821,7 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_src_reg(0, rs1, RegType::Float); instr->set_src_reg(1, rs2, RegType::Float); instr->set_src_reg(2, rs3, RegType::Float); + instr->set_fcsr_use((funct3 == 0x7) ? Instr::FCSR_FRM : 0, Instr::FCSR_FFLAGS); } break; case Opcode::EXT1: { switch (funct7) { @@ -909,7 +933,6 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_op_type(is_sparse ? TcuType::WMMA_SP : TcuType::WMMA); instr->set_args(IntrTcuArgs{0, 0, fmt_s, fmt_d, 0, 0, 0, 0, 0, 0}); instr->set_macro_op(); - instr->set_wstall(true); } break; #ifdef VX_CFG_TCU_WGMMA_ENABLE case 1: { // WGMMA_SYNC — single macro Instr, sequencer expands to micro-ops @@ -920,7 +943,6 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { instr->set_op_type(is_sparse ? TcuType::WGMMA_SP : TcuType::WGMMA); instr->set_args(IntrTcuArgs{is_a_smem ? 1u : 0u, cd_nregs, fmt_s, fmt_d, 0, 0, 0, 0, 0, 0}); instr->set_macro_op(); - instr->set_wstall(true); } break; #endif // VX_CFG_TCU_WGMMA_ENABLE #ifdef TCU_META_ENABLE @@ -955,7 +977,6 @@ Instr::Ptr Decoder::decode(uint32_t code, uint64_t uuid) { std::abort(); } instr->set_macro_op(); - instr->set_wstall(true); // pause fetch while sequencer expands the N uops } break; default: std::abort(); diff --git a/sim/simx/dispatcher.cpp b/sim/simx/dispatcher.cpp index 273c6747cc..f8a45998dd 100644 --- a/sim/simx/dispatcher.cpp +++ b/sim/simx/dispatcher.cpp @@ -16,18 +16,21 @@ using namespace vortex; -Dispatcher::Dispatcher(const SimContext& ctx, const char* name, Core* core, uint32_t buf_size, uint32_t block_size, uint32_t num_lanes) +Dispatcher::Dispatcher(const SimContext& ctx, const char* name, Core* core, uint32_t queue_size, uint32_t block_size, uint32_t num_lanes, uint32_t out_delay) : SimObject(ctx, name) - , Inputs(VX_CFG_ISSUE_WIDTH, this) - // physical block count, matches downstream FU; each output is the per-FU - // dispatch queue, depth = VX_CFG_DISPATCH_QUEUE_SIZE. - , Outputs(block_size, SimChannel(this, buf_size)) + // One dispatch queue per issue slot. The extra entry holds the op still in + // the collector's output register, which the channel counts as occupancy. + , Inputs(VX_CFG_ISSUE_WIDTH, SimChannel(this, queue_size + 1)) + // Bound to the unit's per-block inputs. + , Outputs(block_size, this) + , ReleaseOut(this, VX_CFG_ISSUE_WIDTH) , core_(core) , block_size_(block_size) , num_lanes_(num_lanes) , num_blocks_(VX_CFG_ISSUE_WIDTH / block_size) , num_packets_(VX_CFG_NUM_THREADS / num_lanes) , batch_idx_(0) + , out_delay_(out_delay) , block_pids_(block_size, 0) {} @@ -61,8 +64,8 @@ void Dispatcher::on_tick() { continue; } - // check output buffer capacity — outputs are sized NUM_BLOCKS; - // input[batch_idx*block_size + b] aggregates onto output[b]. + // input[batch_idx*block_size + b] aggregates onto output[b]; the op + // leaves its slot queue only when the unit's input has room. auto& output = Outputs.at(b); if (output.full()) continue; @@ -115,6 +118,7 @@ void Dispatcher::on_tick() { } else { block_pids_.at(b) = -1; // mark block as processed input.pop(); + ReleaseOut.send(trace, 0); ++block_sent; } ThreadMask tmask(VX_CFG_NUM_THREADS); @@ -128,10 +132,11 @@ void Dispatcher::on_tick() { } else { // issue the trace input.pop(); + ReleaseOut.send(trace, 0); ++block_sent; } DT(3, this->name() << "-pipeline dispatch: " << *new_trace); - output.send(new_trace, 1); + output.send(new_trace, out_delay_); } // advance to next batch once all blocks in the current batch have been processed diff --git a/sim/simx/dispatcher.h b/sim/simx/dispatcher.h index 9458dd6b35..3fe1eb9468 100644 --- a/sim/simx/dispatcher.h +++ b/sim/simx/dispatcher.h @@ -25,8 +25,11 @@ class Dispatcher : public SimObject { public: std::vector> Inputs; std::vector> Outputs; + // An op leaving its slot queue, which returns its issue credit. + SimChannel ReleaseOut; - Dispatcher(const SimContext& ctx, const char* name, Core* core, uint32_t buf_size, uint32_t block_size, uint32_t num_lanes); + // out_delay: cycles from leaving the slot queue to the unit taking the op. + Dispatcher(const SimContext& ctx, const char* name, Core* core, uint32_t queue_size, uint32_t block_size, uint32_t num_lanes, uint32_t out_delay); virtual ~Dispatcher(); @@ -41,6 +44,7 @@ class Dispatcher : public SimObject { uint32_t num_blocks_; uint32_t num_packets_; uint32_t batch_idx_; + uint32_t out_delay_; std::vector block_pids_; friend class SimObject; diff --git a/sim/simx/instr.h b/sim/simx/instr.h index 13cc5e1be7..fdb24d2dce 100644 --- a/sim/simx/instr.h +++ b/sim/simx/instr.h @@ -139,6 +139,8 @@ class Instr { , parent_uuid_(uuid) , fu_type_(fu_type) , dst_bytesel_(0xFF) + , fcsr_rd_(0) + , fcsr_wr_(0) , is_uop_(false) , is_macro_op_(false) , is_wstall_(false) @@ -228,6 +230,19 @@ class Instr { void set_dst_bytesel(uint8_t value) { dst_bytesel_ = value; } uint8_t get_dst_bytesel() const { return dst_bytesel_; } + // Floating-point CSR fields the scoreboard orders like GPRs: an instruction + // naming one, read or written, waits for an older writer to commit. + enum : uint8_t { + FCSR_FFLAGS = 1u << 0, + FCSR_FRM = 1u << 1, + }; + void set_fcsr_use(uint8_t rd_mask, uint8_t wr_mask) { + fcsr_rd_ = rd_mask; + fcsr_wr_ = wr_mask; + } + uint8_t fcsr_reads() const { return fcsr_rd_; } + uint8_t fcsr_writes() const { return fcsr_wr_; } + private: uint64_t uuid_; @@ -238,6 +253,8 @@ class Instr { RegOpd rsrc_[MAX_REG_SOURCES]; RegOpd rdest_; uint8_t dst_bytesel_; + uint8_t fcsr_rd_; + uint8_t fcsr_wr_; bool is_uop_; bool is_macro_op_; bool is_wstall_; diff --git a/sim/simx/instr_trace.h b/sim/simx/instr_trace.h index 59c808f4f5..63a275bba8 100644 --- a/sim/simx/instr_trace.h +++ b/sim/simx/instr_trace.h @@ -71,7 +71,8 @@ struct instr_trace_t { // and leave this clear. bool resume_warp; - uint64_t issue_time ; + // Cycle the instruction entered its warp's instruction buffer. + uint64_t ibuf_time; instr_trace_t(uint64_t uuid) : uuid(uuid) @@ -95,7 +96,7 @@ struct instr_trace_t { , num_pkts(1) , fetch_stall(false) , resume_warp(false) - , issue_time(SimPlatform::instance().cycles()) + , ibuf_time(0) , log_once_(false) {} @@ -122,7 +123,7 @@ struct instr_trace_t { , num_pkts(rhs.num_pkts) , fetch_stall(rhs.fetch_stall) , resume_warp(rhs.resume_warp) - , issue_time(rhs.issue_time) + , ibuf_time(rhs.ibuf_time) , log_once_(false) {} diff --git a/sim/simx/mem/lsu_mem_adapter.cpp b/sim/simx/mem/lsu_mem_adapter.cpp index a706d73b87..7e4460c1c5 100644 --- a/sim/simx/mem/lsu_mem_adapter.cpp +++ b/sim/simx/mem/lsu_mem_adapter.cpp @@ -26,6 +26,7 @@ LsuMemAdapter::LsuMemAdapter( , ReqOut(num_inputs, this) , RspIn(num_inputs, this) , delay_(delay) + , port_delay_(num_inputs, 0) , pending_mask_(num_inputs) { assert(num_inputs > 0); @@ -143,7 +144,7 @@ void LsuMemAdapter::on_tick() { out_req.flags.local = (t == AddrType::Shared); } - if (ReqOut.at(i).try_send(out_req, delay_)) { + if (ReqOut.at(i).try_send(out_req, delay_ + port_delay_.at(i))) { DT(4, this->name() << " req" << i << ": " << out_req); pending_mask_.reset(i); } diff --git a/sim/simx/mem/lsu_mem_adapter.h b/sim/simx/mem/lsu_mem_adapter.h index ddbf86579c..0d0e05ebe6 100644 --- a/sim/simx/mem/lsu_mem_adapter.h +++ b/sim/simx/mem/lsu_mem_adapter.h @@ -41,12 +41,16 @@ class LsuMemAdapter : public SimObject { ) : LsuMemAdapter(ctx, name, num_inputs, 0) {} + // Extra request-side registers on one output port. + void set_port_delay(uint32_t port, uint32_t delay) { port_delay_.at(port) = delay; } + protected: void on_reset(); void on_tick(); private: uint32_t delay_; + std::vector port_delay_; BitVector pending_mask_; friend class SimObject; diff --git a/sim/simx/mem/mem_coalescer.cpp b/sim/simx/mem/mem_coalescer.cpp index 964177dd50..533e236b13 100644 --- a/sim/simx/mem/mem_coalescer.cpp +++ b/sim/simx/mem/mem_coalescer.cpp @@ -39,6 +39,7 @@ MemCoalescer::MemCoalescer( , sent_mask_(input_size) , line_size_(line_size) , delay_(delay) + , port_delay_(output_size, 0) {} void MemCoalescer::on_reset() { @@ -294,7 +295,7 @@ void MemCoalescer::flush_out_round() { if (!out_round_.lanes.test(o)) { continue; } - if (ReqOut.at(o).try_send(out_round_.reqs.at(o), delay_)) { + if (ReqOut.at(o).try_send(out_round_.reqs.at(o), delay_ + port_delay_.at(o))) { out_round_.lanes.reset(o); } } diff --git a/sim/simx/mem/mem_coalescer.h b/sim/simx/mem/mem_coalescer.h index 4db68c5e50..e133c4be1a 100644 --- a/sim/simx/mem/mem_coalescer.h +++ b/sim/simx/mem/mem_coalescer.h @@ -46,6 +46,9 @@ class MemCoalescer : public SimObject { const PerfStats& perf_stats() const; + // Extra request-side registers on one output port. + void set_port_delay(uint32_t port, uint32_t delay) { port_delay_.at(port) = delay; } + protected: void on_reset(); void on_tick(); @@ -76,6 +79,7 @@ class MemCoalescer : public SimObject { out_round_t out_round_; uint32_t line_size_; uint32_t delay_; + std::vector port_delay_; PerfStats perf_stats_; friend class SimObject; diff --git a/sim/simx/om/om_unit.cpp b/sim/simx/om/om_unit.cpp index 5d6a719c2a..00f737de39 100644 --- a/sim/simx/om/om_unit.cpp +++ b/sim/simx/om/om_unit.cpp @@ -53,6 +53,9 @@ Instr::Ptr OmUopGen::get(const Instr& macro_instr, uint32_t uop_index) { IntrOmArgs uopArgs{}; uopArgs.export_mask = (uop_index + 1 == total) ? (args.export_mask & 0x3) : 0; uop_instr->set_args(uopArgs); + // No other warp of the issue slot issues between the two beats of a record. + uop_instr->set_fu_lock(uop_index == 0); + uop_instr->set_fu_unlock(uop_index + 1 == total); return uop_instr; } diff --git a/sim/simx/rtu/rtu_unit.cpp b/sim/simx/rtu/rtu_unit.cpp index 37799aaf36..5b6a93ad8a 100644 --- a/sim/simx/rtu/rtu_unit.cpp +++ b/sim/simx/rtu/rtu_unit.cpp @@ -394,6 +394,10 @@ Instr::Ptr RtuUopGen::get(const Instr& macro_instr, uint32_t uop_index) { default: std::abort(); } + // The RTU takes one ray at a time: no other warp of the issue slot issues + // between its beats. + uop->set_fu_lock(uop_index == 0); + uop->set_fu_unlock(uop_index == total - 1); } else { std::abort(); // only TRACE / GETWF / GETW are SFU macro-ops } diff --git a/sim/simx/scoreboard.cpp b/sim/simx/scoreboard.cpp index fee0fece6d..7701aa15cb 100644 --- a/sim/simx/scoreboard.cpp +++ b/sim/simx/scoreboard.cpp @@ -12,12 +12,14 @@ // limitations under the License. #include "scoreboard.h" +#include using namespace vortex; Scoreboard::Scoreboard(const SimContext& ctx, const char* name) : SimObject(ctx, name) - , in_use_regs_(VX_CFG_NUM_WARPS) { + , in_use_regs_(VX_CFG_NUM_WARPS) + , in_use_fcsr_(VX_CFG_NUM_WARPS, 0) { for (auto& in_use_reg : in_use_regs_) { in_use_reg.resize((int)RegType::Count); } @@ -33,6 +35,7 @@ void Scoreboard::on_reset() { mask.reset(); } } + std::fill(in_use_fcsr_.begin(), in_use_fcsr_.end(), 0); owners_.clear(); commit_counts_.clear(); pending_releases_.clear(); @@ -42,8 +45,12 @@ void Scoreboard::on_tick() { uint64_t now = SimPlatform::instance().cycles(); while (!pending_releases_.empty() && pending_releases_.front().due <= now) { auto& r = pending_releases_.front(); - owners_.erase(get_reg_id(r.reg, r.wid)); - in_use_regs_.at(r.wid).at((int)r.reg.type).reset(r.reg.idx); + if (r.fcsr != 0) { + in_use_fcsr_.at(r.wid) &= ~r.fcsr; + } else { + owners_.erase(get_reg_id(r.reg, r.wid)); + in_use_regs_.at(r.wid).at((int)r.reg.type).reset(r.reg.idx); + } pending_releases_.pop_front(); } if (pending_releases_.empty()) { @@ -52,6 +59,10 @@ void Scoreboard::on_tick() { } bool Scoreboard::in_use(instr_trace_t* trace) const { + auto& instr = *trace->instr_ptr; + if (in_use_fcsr_.at(trace->wid) & (instr.fcsr_reads() | instr.fcsr_writes())) { + return true; + } if (trace->wb) { assert(trace->dst_reg.type != RegType::None); if (in_use_regs_.at(trace->wid).at((int)trace->dst_reg.type).test(trace->dst_reg.idx)) { @@ -104,7 +115,21 @@ void Scoreboard::release(instr_trace_t* trace) { assert(in_use_regs_.at(trace->wid).at((int)trace->dst_reg.type).test(trace->dst_reg.idx)); assert(owners_.count(reg_id) != 0); commit_counts_.erase(reg_id); - pending_releases_.push_back({trace->wid, trace->dst_reg, + pending_releases_.push_back({trace->wid, trace->dst_reg, 0, + SimPlatform::instance().cycles() + kReleaseDelay}); + this->tick_wake(); +} + +void Scoreboard::reserve_fcsr(instr_trace_t* trace) { + uint8_t fields = trace->instr_ptr->fcsr_writes(); + assert((in_use_fcsr_.at(trace->wid) & fields) == 0); + in_use_fcsr_.at(trace->wid) |= fields; +} + +void Scoreboard::release_fcsr(instr_trace_t* trace) { + uint8_t fields = trace->instr_ptr->fcsr_writes(); + assert((in_use_fcsr_.at(trace->wid) & fields) == fields); + pending_releases_.push_back({trace->wid, {RegType::None, 0}, fields, SimPlatform::instance().cycles() + kReleaseDelay}); this->tick_wake(); } diff --git a/sim/simx/scoreboard.h b/sim/simx/scoreboard.h index 2cad40ea5b..e05ade6479 100644 --- a/sim/simx/scoreboard.h +++ b/sim/simx/scoreboard.h @@ -54,6 +54,12 @@ class Scoreboard : public SimObject { // but cache responses may complete out of order, so we count COMMITS. bool commit_packet(instr_trace_t* trace); + // FCSR fields (Instr::FCSR_*) the instruction writes: claimed at + // issue, freed on the same schedule as a destination register once the + // instruction's last packet commits. + void reserve_fcsr(instr_trace_t* trace); + void release_fcsr(instr_trace_t* trace); + protected: void on_reset(); void on_tick(); @@ -63,15 +69,17 @@ class Scoreboard : public SimObject { return (wid << RegOpd::ID_BITS) | reg.id(); } - static constexpr uint32_t kReleaseDelay = 2; + static constexpr uint32_t kReleaseDelay = 1; struct pending_release_t { uint32_t wid; RegOpd reg; + uint8_t fcsr; uint64_t due; }; std::vector> in_use_regs_; + std::vector in_use_fcsr_; std::deque pending_releases_; std::unordered_map owners_; std::unordered_map commit_counts_; diff --git a/sim/simx/sequencer.h b/sim/simx/sequencer.h index 558d83a1a4..9f7222f270 100644 --- a/sim/simx/sequencer.h +++ b/sim/simx/sequencer.h @@ -48,6 +48,9 @@ class Sequencer : public SimObject { // Advance to next micro-op. Returns true when all micro-ops have been issued. bool advance(); + // A macro-op's first micro-op is the current one. + bool starting() const { return state_.active && state_.uop_index == 0; } + // Drop any cached uop / macro-op state for this sequencer. Used by // Core::flush_warp_pipeline at async-trap entry: when the ibuffer is // flushed, the trace pointer this sequencer cached in state_.current_uop diff --git a/sim/simx/tcu/tcu_unit.cpp b/sim/simx/tcu/tcu_unit.cpp index e31b65ac68..5e2bd89173 100644 --- a/sim/simx/tcu/tcu_unit.cpp +++ b/sim/simx/tcu/tcu_unit.cpp @@ -393,6 +393,9 @@ class TcuUnit::Impl { } #endif exec_done_.fill(false); + for (auto& due : result_due_) { + due.clear(); + } in_wgmma_.fill(false); lmem_desc_.clear(); wgmma_desc_.fill({0, 0}); @@ -560,10 +563,21 @@ class TcuUnit::Impl { uint32_t setup_fired = 0; #endif + uint64_t now = SimPlatform::instance().cycles(); for (uint32_t b = 0; b < VX_CFG_NUM_TCU_BLOCKS; ++b) { + // A result completing while an earlier one still waits for the commit + // side holds admission for the cycle. Results already visible but not + // taken are the output occupancy beyond those still in flight. + auto& due = result_due_.at(b); + due.erase(std::remove_if(due.begin(), due.end(), [now](uint64_t t) { return t < now; }), due.end()); + bool completing = std::any_of(due.begin(), due.end(), [now](uint64_t t) { return t == now; }); + bool backed_up = simobject_->Outputs.at(b).size() > due.size(); auto& input = simobject_->Inputs.at(b); if (input.empty()) continue; + if (completing && backed_up) { + continue; + } auto trace = input.peek(); auto tcu_type = std::get(trace->op_type); auto tpuArgs = std::get(trace->instr_ptr->get_args()); @@ -682,6 +696,7 @@ class TcuUnit::Impl { } #endif if (simobject_->Outputs.at(b).try_send(trace, delay)) { + due.push_back(now + delay); exec_done_.at(b) = false; #ifdef TCU_META_ENABLE // TCU_LD retired: free the AGU for the next metadata load. @@ -1549,6 +1564,8 @@ class TcuUnit::Impl { mutable PerfStats perf_stats_; // Per-block guard: execute already happened for this trace; reset on pop(). std::array exec_done_; + // Cycles at which each block's results sent so far become visible. + std::array, VX_CFG_NUM_TCU_BLOCKS> result_due_; // True while a block is between its first and last WGMMA uop. std::array in_wgmma_; // Current block index, set before delegating to wgmma(). From aa23cad96f74ca17b5228df2efcfd133ff54a001 Mon Sep 17 00:00:00 2001 From: tinebp Date: Sat, 3 Oct 2026 04:00:47 -0700 Subject: [PATCH 8/8] ci: run the synthesis gates only on nights that follow a synthesis-input change GitHub cannot condition a cron schedule on the repository. Both gates therefore started a run every night just to find out whether master had moved. On a quiet night that left an empty run behind. And because the check keyed on master's SHA, a docs-only or SimX-only commit triggered a full multi-hour sweep. ASIC gate (hosted runners, no host to keep state on) now arms and disarms itself: - The workflow stays disabled between qualifying pushes, so its schedule does not fire. - ci.yml's plan job enables it on a push to master that touches hw/, the configuration TOMLs, VERSION, the gate's scripts, catalog or baselines, or the workflow itself. An indeterminate diff also arms it. - The scheduled run disables the workflow as its first step, so a push that lands mid-sweep re-arms the next night. A build error re-enables it to retry. A pass or a regression is a verdict, and its run is the record. - The SHA-marker cache and the skipped-night republish are gone. FPGA gate (self-hosted Vivado runner) loses its schedule: - A nightly cron on the runner's host runs ci/fpga_gate_dispatch.sh. It dispatches the workflow only when a synthesis input changed since the last gated SHA (the workflow's existing state file), the runner is up, and no gate run is queued or in progress. - The workflow drops its skip check, its republish step and the force input. A dispatched run always has work. synth_report.py loses --note: only the removed republish steps used it. Host setup still needed: install the crontab line from the script's header, and authenticate gh with a token that can dispatch workflows. Not run: these workflows only execute on GitHub. The dispatcher passed bash -n but has not been exercised end to end. Co-Authored-By: Claude Opus 5.5 (1M context) --- .github/workflows/asic_gate.yml | 119 ++++++++----------------- .github/workflows/ci.yml | 23 +++++ .github/workflows/fpga_gate.yml | 91 +++++-------------- ci/fpga_gate_dispatch.sh | 55 ++++++++++++ ci/synth_report.py | 8 +- docs/designs/continuous_integration.md | 40 ++++++--- 6 files changed, 167 insertions(+), 169 deletions(-) create mode 100755 ci/fpga_gate_dispatch.sh diff --git a/.github/workflows/asic_gate.yml b/.github/workflows/asic_gate.yml index febd0ca2ba..a570a460b1 100644 --- a/.github/workflows/asic_gate.yml +++ b/.github/workflows/asic_gate.yml @@ -13,10 +13,13 @@ # - A DUT takes 1-2 hours, so the builds must fan out to ONE STANDALONE JOB # EACH. A ci.yml cell is one job running a pytest slice; eight sequential # synthesis runs in one cell would be a day. -# - Nightly builds are expensive even when free, so the run is SKIPPED unless -# master has actually moved since the last gate run. Unlike the self-hosted -# fpga_gate, a hosted runner keeps no state between runs, so the "already -# gated this commit" marker lives in the actions cache, keyed by SHA. +# - Nightly builds are expensive even when free, so the nightly runs only +# after a push changed what it synthesizes. GitHub cannot condition a +# schedule on the repository, so the workflow stays DISABLED between those +# pushes: ci.yml's plan job enables it on a push to master that touches the +# synthesis inputs, and the scheduled run disables it again as it starts. +# A quiet night therefore creates no run at all. A manual run needs the +# workflow enabled first (`gh workflow enable asic_gate.yml`). # # It always gates master, whatever branch the schedule happens to fire on. No # licence and no dedicated machine: yosys/sv2v/OpenSTA ship in the prebuilt @@ -34,13 +37,9 @@ on: builds: description: "build ids/groups, space separated (blank = all)" default: "" - force: - description: "run even if master has not moved since the last gate" - type: boolean - default: false -# One sweep at a time, so two nights' runs cannot both claim the SHA marker. -# Never cancel a run in flight -- hours of synthesis are already spent. +# One sweep at a time. Never cancel a run in flight -- hours of synthesis are +# already spent. concurrency: group: asic-gate cancel-in-progress: false @@ -50,77 +49,43 @@ env: jobs: # --------------------------------------------------------------------------- - # plan — decide whether to run at all, and emit one matrix entry per build. + # plan — disarm the nightly, and emit one matrix entry per build. # No build env; just PyYAML. # --------------------------------------------------------------------------- plan: runs-on: ubuntu-22.04 + permissions: + actions: write + contents: read outputs: run: ${{ steps.q.outputs.run }} builds: ${{ steps.q.outputs.builds }} sha: ${{ steps.q.outputs.sha }} - key: ${{ steps.q.outputs.key }} steps: + # First, so a push that lands while this run synthesizes re-arms the next + # night instead of being cleared when this one ends. A failure here only + # costs an extra night, so it warns rather than blocking the gate. + - name: Disarm until the next push + if: github.event_name == 'schedule' + env: + GH_TOKEN: ${{ github.token }} + run: | + gh api -X PUT "repos/$GITHUB_REPOSITORY/actions/workflows/asic_gate.yml/disable" \ + || echo "::warning::could not disable asic_gate.yml; it will run again tomorrow" + - uses: actions/checkout@v4 with: ref: master # hard-pinned: the gate always tracks master - fetch-depth: 0 - run: pip install --quiet pyyaml - # The marker key is master's head AS CHECKED OUT, not github.sha: on a - # schedule those are normally the same, but the thing being gated is what - # the checkout produced. Computed here rather than with hashFiles() in the - # step below, which cannot see a runtime value. - - name: Compute gate key - id: key - run: | - set -euo pipefail - SHA=$(git rev-parse HEAD) - # The spec is in the key too, so editing the build list re-gates a - # commit that was already gated under the old list. - SPEC=$(sha256sum ci/testcases/asic_gate.yaml | cut -c1-16) - echo "sha=$SHA" >> "$GITHUB_OUTPUT" - echo "key=asic-gate-$SHA-$SPEC" >> "$GITHUB_OUTPUT" - - # Restore this commit's marker. A match means a previous run already - # reached a verdict on this SHA, so master has not moved and there is - # nothing new to gate. The `report` job writes the marker at the end, with - # that run's reports inside it for a skipped night to republish. A cache - # entry cannot be overwritten, so each run saves under its own key and the - # newest one for the commit is what restores. - - name: Read gate marker - id: marker - uses: actions/cache/restore@v4 - with: - path: .asic_gate - key: ${{ steps.key.outputs.key }} - restore-keys: ${{ steps.key.outputs.key }}- - - name: Plan builds id: q env: IN_BUILDS: ${{ github.event.inputs.builds }} - FORCE: ${{ github.event.inputs.force }} - HIT: ${{ steps.marker.outputs.cache-matched-key != '' }} - SHA: ${{ steps.key.outputs.sha }} - KEY: ${{ steps.key.outputs.key }} - RUNS: ${{ github.server_url }}/${{ github.repository }}/actions/workflows/asic_gate.yml run: | set -euo pipefail + SHA=$(git rev-parse HEAD) echo "sha=$SHA" >> "$GITHUB_OUTPUT" - echo "key=$KEY" >> "$GITHUB_OUTPUT" - if [ "$HIT" = "true" ] && [ "${FORCE:-}" != "true" ]; then - echo "master unchanged since the last gate run ($SHA) — skipping" - echo "run=false" >> "$GITHUB_OUTPUT" - echo "builds=[]" >> "$GITHUB_OUTPUT" - # A skipped night has no verdict of its own, so it republishes the - # one the gated commit got: an unchanged red master must not read - # as an empty run. - NOTE="master is unchanged since the last gate run (\`$SHA\`), so nothing was synthesized. These are the results of [that run]($(cat .asic_gate/url 2>/dev/null || echo "$RUNS"))." - python3 ci/synth_report.py --annotate warning --note "$NOTE" .asic_gate \ - >> "$GITHUB_STEP_SUMMARY" || true - exit 0 - fi echo "gating $SHA" # Dispatch inputs only narrow the catalog, and reach the shell through # an env var: interpolating one into the script would let a dispatch @@ -241,13 +206,16 @@ jobs: esac # --------------------------------------------------------------------------- - # report — collect every build's verdict into one summary, and record the SHA - # so tomorrow's run skips an unchanged master. + # report — collect every build's verdict into one summary, and re-arm the + # nightly when a build could not reach one. # --------------------------------------------------------------------------- report: needs: [plan, synth] if: always() && needs.plan.outputs.run == 'true' runs-on: ubuntu-22.04 + permissions: + actions: write + contents: read steps: - uses: actions/checkout@v4 with: @@ -263,7 +231,7 @@ jobs: # replacing it. Without the guard, synth_report.py's non-zero exit (1 = # regression, 2 = build error -- i.e. every case this step exists to # report) kills the step on that line, so `rc` never reaches - # $GITHUB_OUTPUT and the marker steps below read it as "". + # $GITHUB_OUTPUT and the re-arm step below reads it as "". - name: Summarize id: sum run: | @@ -276,25 +244,16 @@ jobs: echo "rc=$rc" >> "$GITHUB_OUTPUT" exit 0 - # Record the SHA once every build has reached a VERDICT (pass or - # regression), so a red master is not re-synthesized every night — the - # failed run is the record. A build error (rc=2) does NOT record, so the - # next nightly retries it. - - name: Record gated SHA - if: steps.sum.outputs.rc != '2' + # A pass or a regression is a verdict, and the run is its record: a red + # master is not re-synthesized every night. A build error (rc=2) is not, + # so it re-arms the nightly to retry without waiting for another push. + - name: Re-arm after a build error + if: steps.sum.outputs.rc == '2' env: - RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }} + GH_TOKEN: ${{ github.token }} run: | - mkdir -p .asic_gate - echo "${{ needs.plan.outputs.sha }}" > .asic_gate/sha - echo "$RUN_URL" > .asic_gate/url - cp reports/asic_gate_*.json .asic_gate/ - - name: Save gate marker - if: steps.sum.outputs.rc != '2' - uses: actions/cache/save@v4 - with: - path: .asic_gate - key: ${{ needs.plan.outputs.key }}-${{ github.run_id }}-${{ github.run_attempt }} + gh api -X PUT "repos/$GITHUB_REPOSITORY/actions/workflows/asic_gate.yml/enable" \ + || echo "::warning::could not re-enable asic_gate.yml; the next push re-arms it" # Unzipped, so the summary opens in the browser straight from the run # page, and the joined report it was rendered from downloads as itself. diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f915dd7e35..d8bdfcc7e0 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -33,6 +33,9 @@ jobs: # --------------------------------------------------------------------------- plan: runs-on: ubuntu-22.04 + permissions: + actions: write + contents: read outputs: cells: ${{ steps.q.outputs.cells }} run: ${{ steps.q.outputs.run }} @@ -41,6 +44,26 @@ jobs: - uses: actions/checkout@v4 with: fetch-depth: 0 + # The ASIC gate's nightly stays disabled until a push gives it something + # to synthesize; it disables itself again when it starts. An + # indeterminate diff arms it. Arming never fails CI: a miss only delays + # the gate to the next qualifying push. + - name: Arm the ASIC gate + if: github.event_name == 'push' && github.ref == 'refs/heads/master' + env: + GH_TOKEN: ${{ github.token }} + BASE: ${{ github.event.before }} + run: | + set -uo pipefail + INPUTS='^(hw/|VX_config\.toml$|VX_types\.toml$|VERSION$|ci/(asic_gate|synth_gate|synth_report)\.py$|ci/testcases/asic_gate\.yaml$|ci/baselines/synthesis/yosys/|\.github/workflows/asic_gate\.yml$)' + if git cat-file -e "${BASE}^{commit}" 2>/dev/null \ + && ! git diff --name-only "$BASE" HEAD | grep -qE "$INPUTS"; then + echo "no synthesis input changed; the ASIC gate stays as it is" + exit 0 + fi + gh api -X PUT "repos/$GITHUB_REPOSITORY/actions/workflows/asic_gate.yml/enable" \ + && echo "ASIC gate armed for tonight" \ + || echo "::warning::could not enable asic_gate.yml; the next qualifying push retries" - run: pip install --quiet pyyaml # Read the toolchain pin so we can probe its cache below. TOOLCHAIN_REV is # the toolchain cache key; the build job's setup-vortex caches `tools` under diff --git a/.github/workflows/fpga_gate.yml b/.github/workflows/fpga_gate.yml index 6bb7b57cf5..ab154187a9 100644 --- a/.github/workflows/fpga_gate.yml +++ b/.github/workflows/fpga_gate.yml @@ -2,22 +2,21 @@ # # Licensed under the Apache License, Version 2.0 (the "License"). # -# fpga_gate — nightly FPGA synthesis-regression gate on the self-hosted Vivado -# runner. Synthesizes the DUT catalog (ci/testcases/fpga_gate.yaml) and asserts -# Fmax/LUT against the golden results in ci/baselines/synthesis/xilinx/, so an -# RTL change that costs timing closure or area cannot sit unnoticed on master. +# fpga_gate — FPGA synthesis-regression gate on the self-hosted Vivado runner. +# Synthesizes the DUT catalog (ci/testcases/fpga_gate.yaml) and asserts Fmax/LUT +# against the golden results in ci/baselines/synthesis/xilinx/, so an RTL change +# that costs timing closure or area cannot sit unnoticed on master. # -# Nightly builds are hours long, so the run is SKIPPED unless origin/master has -# actually moved since the last gate run on this runner (state file below). It -# always gates origin/master, whatever branch the schedule happens to fire on. +# Builds are hours long, so the workflow has no schedule of its own: a nightly +# cron on the runner's host runs ci/fpga_gate_dispatch.sh, which dispatches it +# only when a synthesis input changed since the last gated commit (the state +# file below). A quiet night creates no run. It always gates origin/master. # # See docs/designs/continuous_integration.md §3.5. name: FPGA Gate on: - schedule: - - cron: '0 3 * * *' # nightly 03:00 UTC workflow_dispatch: inputs: builds: @@ -26,10 +25,6 @@ on: jobs: description: "max parallel Vivado builds" default: "2" - force: - description: "run even if origin/master has not moved" - type: boolean - default: false # One synthesis sweep at a time: parallel Vivado runs would oversubscribe the # machine and make Fmax non-comparable. Never cancel a run in flight. @@ -48,48 +43,15 @@ jobs: fetch-depth: 0 submodules: recursive - # Vivado builds are hours long; only pay for them when master moved. - - name: Check master for new commits + # The state file is the dispatcher's record of the last gated commit. + - name: Resolve the gated commit id: gate - env: - FORCE: ${{ github.event.inputs.force }} run: | set -euo pipefail SHA=$(git rev-parse HEAD) - STATE="$HOME/.cache/vortex/fpga_gate.$(echo '${{ github.repository }}' | tr / _).sha" - LAST=$(cat "$STATE" 2>/dev/null || echo "") - echo "sha=$SHA" >> "$GITHUB_OUTPUT" - echo "state=$STATE" >> "$GITHUB_OUTPUT" - echo "report=${STATE%.sha}.report.json" >> "$GITHUB_OUTPUT" - echo "url=${STATE%.sha}.url" >> "$GITHUB_OUTPUT" - if [ "$SHA" = "$LAST" ] && [ "$FORCE" != "true" ]; then - echo "master unchanged since the last gate run ($SHA) — skipping" - echo "run=false" >> "$GITHUB_OUTPUT" - else - echo "gating $SHA (last run: ${LAST:-none})" - echo "run=true" >> "$GITHUB_OUTPUT" - fi - - # A skipped night has no verdict of its own, so it republishes the one the - # gated commit got: an unchanged red master must not read as an empty run. - - name: Republish the gated verdict - if: steps.gate.outputs.run != 'true' - env: - SHA: ${{ steps.gate.outputs.sha }} - REPORT: ${{ steps.gate.outputs.report }} - URL: ${{ steps.gate.outputs.url }} - RUNS: ${{ github.server_url }}/${{ github.repository }}/actions/workflows/fpga_gate.yml - run: | - set -uo pipefail - NOTE="master is unchanged since the last gate run (\`$SHA\`), so nothing was synthesized." - if [ -f "$REPORT" ]; then - NOTE="$NOTE These are the results of [that run]($(cat "$URL" 2>/dev/null || echo "$RUNS"))." - python3 ci/synth_report.py --annotate warning --note "$NOTE" "$REPORT" \ - >> "$GITHUB_STEP_SUMMARY" || true - else - printf '## xilinx gate\n\n%s This runner holds no saved results for it; see the [earlier runs](%s).\n' \ - "$NOTE" "$RUNS" >> "$GITHUB_STEP_SUMMARY" - fi + echo "sha=$SHA" >> "$GITHUB_OUTPUT" + echo "state=$HOME/.cache/vortex/fpga_gate.$(echo '${{ github.repository }}' | tr / _).sha" >> "$GITHUB_OUTPUT" + echo "gating $SHA" # The command is the one the catalog declares (ci/testcases/fpga_gate.yaml's # single `run:` case), so what CI runs and what the catalog says cannot @@ -98,7 +60,6 @@ jobs: # execute arbitrary commands on the runner. - name: Run fpga_gate id: run - if: steps.gate.outputs.run == 'true' env: IN_BUILDS: ${{ github.event.inputs.builds }} IN_JOBS: ${{ github.event.inputs.jobs }} @@ -118,7 +79,7 @@ jobs: # not pass did not, on the run's own page. The same document is uploaded # below as a file a browser opens directly. - name: Summarize - if: always() && steps.gate.outputs.run == 'true' + if: always() run: | set -uo pipefail python3 ci/synth_report.py --annotate error fpga_gate_report.json \ @@ -130,7 +91,7 @@ jobs: # The log is printed with workflow commands stopped, so nothing in it can # be read as one. - name: Failed build logs - if: always() && steps.gate.outputs.run == 'true' + if: always() run: | set -uo pipefail TOKEN="log-$(head -c 16 /dev/urandom | od -An -tx1 | tr -d ' \n')" @@ -148,25 +109,15 @@ jobs: # Record the SHA once the gate has actually reached a verdict (pass=0 or # regression=1), so a red master is not re-synthesized every night — the # failed run is the record. An infra/build failure (2) does NOT record, so - # the next nightly retries it. The report and the run's URL are kept next - # to it, for the skipped nights to republish. + # the next night's dispatcher retries it. - name: Record gated SHA - if: steps.gate.outputs.run == 'true' && steps.run.outputs.rc != '2' - env: - REPORT: ${{ steps.gate.outputs.report }} - URL: ${{ steps.gate.outputs.url }} - RUN_URL: ${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }} + if: steps.run.outputs.rc == '0' || steps.run.outputs.rc == '1' run: | mkdir -p "$(dirname '${{ steps.gate.outputs.state }}')" echo "${{ steps.gate.outputs.sha }}" > "${{ steps.gate.outputs.state }}" - rm -f "$REPORT" "$URL" - if [ -f fpga_gate_report.json ]; then - cp fpga_gate_report.json "$REPORT" - echo "$RUN_URL" > "$URL" - fi - uses: actions/upload-artifact@v4 - if: always() && steps.gate.outputs.run == 'true' + if: always() with: name: fpga-gate-report path: | @@ -182,14 +133,14 @@ jobs: # page, and the report it was rendered from downloads as itself. The # report holds everything a baseline records. - uses: actions/upload-artifact@v7 - if: always() && steps.gate.outputs.run == 'true' + if: always() with: path: fpga_gate_report.md archive: false overwrite: true if-no-files-found: warn - uses: actions/upload-artifact@v7 - if: always() && steps.gate.outputs.run == 'true' + if: always() with: path: fpga_gate_report.json archive: false @@ -198,7 +149,7 @@ jobs: # always(), not success(): the verdict must not be masked by an upload failure. - name: Verdict - if: always() && steps.gate.outputs.run == 'true' + if: always() run: | case "${{ steps.run.outputs.rc }}" in 0) echo "fpga_gate passed" ;; diff --git a/ci/fpga_gate_dispatch.sh b/ci/fpga_gate_dispatch.sh new file mode 100755 index 0000000000..63811ce3d2 --- /dev/null +++ b/ci/fpga_gate_dispatch.sh @@ -0,0 +1,55 @@ +#!/bin/bash +# Nightly dispatcher for the FPGA gate (.github/workflows/fpga_gate.yml), run +# from cron on the host of the self-hosted Vivado runner: +# +# 0 3 * * * /ci/fpga_gate_dispatch.sh >> ~/.cache/vortex/fpga_gate_dispatch.log 2>&1 +# +# GitHub cannot condition a schedule on the repository, so the workflow has no +# schedule of its own. This dispatches it only when a synthesis input changed +# since the last gated commit (the state file the workflow records), so a quiet +# night creates no run. `gh` must be authenticated with a token that can +# dispatch workflows on the repository. +# +# See docs/designs/continuous_integration.md §4.4. + +set -euo pipefail + +REPO="${REPO:-vortexgpgpu/vortex}" +CLONE="$(cd "$(dirname "$0")/.." && pwd)" +STATE="$HOME/.cache/vortex/fpga_gate.${REPO//\//_}.sha" +INPUTS='^(hw/|VX_config\.toml$|VX_types\.toml$|VERSION$|ci/(fpga_gate|synth_gate|synth_report)\.py$|ci/testcases/fpga_gate\.yaml$|ci/baselines/synthesis/xilinx/|\.github/workflows/fpga_gate\.yml$)' + +log() { + echo "$(date -u +%FT%TZ) fpga_gate_dispatch: $*" +} + +# A run nobody can pick up waits in the queue until GitHub cancels it. +if ! pgrep -x Runner.Listener > /dev/null; then + log "the Actions runner is not running on $(hostname); nothing dispatched" + exit 1 +fi + +# The concurrency group would serialize a second run, but it would still sit +# in the history as a duplicate. +for status in queued in_progress; do + n=$(gh api "repos/$REPO/actions/workflows/fpga_gate.yml/runs?branch=master&status=$status" --jq '.total_count') + if [ "$n" != "0" ]; then + log "a gate run is already $status; nothing dispatched" + exit 0 + fi +done + +git -C "$CLONE" fetch --quiet origin master +TIP=$(git -C "$CLONE" rev-parse FETCH_HEAD) +LAST=$(cat "$STATE" 2>/dev/null || true) + +# A missing or unknown last commit has no diff to read, so it gates. +if [ -n "$LAST" ] && git -C "$CLONE" cat-file -e "${LAST}^{commit}" 2>/dev/null; then + if ! git -C "$CLONE" diff --name-only "$LAST" "$TIP" | grep -qE "$INPUTS"; then + log "no synthesis input changed since $LAST; nothing dispatched" + exit 0 + fi +fi + +gh api -X POST "repos/$REPO/actions/workflows/fpga_gate.yml/dispatches" -f ref=master +log "dispatched the gate for $TIP (last gated: ${LAST:-none})" diff --git a/ci/synth_report.py b/ci/synth_report.py index 73db163630..03356ac238 100755 --- a/ci/synth_report.py +++ b/ci/synth_report.py @@ -13,8 +13,7 @@ 1 at least one build regressed 2 at least one build never produced metrics -- the commit is NOT gated - usage: synth_report.py [--note TEXT] [--annotate LEVEL] [--json FILE] - [--build-failed] + usage: synth_report.py [--annotate LEVEL] [--json FILE] [--build-failed] [...] """ @@ -201,9 +200,6 @@ def main(argv=None): description="Render synthesis-gate reports as one Markdown summary") ap.add_argument("reports", nargs="+", metavar="PATH", help="report JSON, or a directory searched for them") - ap.add_argument("--note", metavar="TEXT", - help="a line printed under the heading, e.g. which run " - "these results came from") ap.add_argument("--annotate", choices=("error", "warning", "notice"), metavar="LEVEL", help="emit a workflow annotation (error, warning or notice) " @@ -235,8 +231,6 @@ def main(argv=None): return 0 print("## %s gate\n" % (tool or "synthesis")) - if args.note: - print("%s\n" % args.note) if not builds: print("No reports found — every build errored before it could write " "one.") diff --git a/docs/designs/continuous_integration.md b/docs/designs/continuous_integration.md index 7bc048fd76..04238ea748 100644 --- a/docs/designs/continuous_integration.md +++ b/docs/designs/continuous_integration.md @@ -369,7 +369,7 @@ number into a location. They are never gated. | job | does | |---|---| -| `plan` | applies the event policy and the escalations, runs `testcase.py matrix` and `lint`, emits the cells | +| `plan` | applies the event policy and the escalations, runs `testcase.py matrix` and `lint`, emits the cells; on a push to master that touches a synthesis input, arms the ASIC gate (§4.5) | | `setup` | warms the toolchain and third-party caches once; a no-op on a cache hit | | `build` | one build tree per XLEN, uploaded as an artifact | | `tests` | one job per cell: `pytest ci -m ""`, JUnit out | @@ -421,18 +421,25 @@ container. ### 4.4 `fpga_gate.yml` -Nightly, on the self-hosted runner with Vivado. It **hard-pins -`origin/master`**: whatever branch the schedule fires on, what is gated is -master's head. +On the self-hosted runner with Vivado. It **hard-pins `origin/master`**: +what is gated is master's head. + +GitHub cannot condition a schedule on the repository: a cron workflow that +checks for changes leaves an empty run behind on every quiet night. So the +workflow has **no schedule**. A nightly cron on the runner's host runs +[`ci/fpga_gate_dispatch.sh`](../../ci/fpga_gate_dispatch.sh), which +dispatches the workflow only when a synthesis input changed since the last +gated commit, and only when the runner is up and no gate run is queued or in +progress. | behavior | detail | |---|---| -| skip when unchanged | the runner keeps the last gated SHA; a run whose master matches is a no-op unless forced | -| when the SHA is recorded | once the gate reaches a verdict, pass or regression. A red master is not re-synthesized every night — the failed run is the record. A build error does not record, so the next night retries | +| when it runs | the dispatcher found a change under `hw/`, the configuration TOMLs, `VERSION`, the gate's scripts, catalog or baselines, or the workflow itself; or a manual dispatch | +| when the SHA is recorded | once the gate reaches a verdict, pass or regression. A red master is not re-synthesized every night — the failed run is the record. A build error does not record, so the next night's dispatcher retries | +| host setup | the crontab line in the script's header, and `gh` authenticated with a token that can dispatch workflows | | diagnosis | the workflow summary is [`synth_report.py`](../../ci/synth_report.py)'s rendering of the report: every build's metrics against its baseline, the reason behind each verdict that is not a pass, and the critical paths and high-fanout nets | | annotations | one per finding, naming the build and the metric | | artifacts | the summary and the report JSON as files a browser opens directly; every build's log and reports as one archive | -| a skipped night | republishes the gated commit's results with a link to the run that produced them, so an unchanged red master does not read as an empty run | ### 4.5 `asic_gate.yml` @@ -442,14 +449,23 @@ out to one job each**. | job | does | |---|---| -| `plan` | reads the spec and emits one matrix entry per build, longest first | +| `plan` | disarms the nightly, reads the spec and emits one matrix entry per build, longest first | | `toolchain` | warms the toolchain cache ahead of the matrix | | `synth` | one build per job, `fail-fast: false` — each design is an independent measurement | -| `report` | joins the per-job reports into one summary and records the SHA marker | +| `report` | joins the per-job reports into one summary; re-arms the nightly after a build error | + +It pins master. A hosted runner has no host to run a dispatcher from, so the +nightly **arms and disarms itself**. The workflow stays disabled between +qualifying pushes, and a disabled workflow's schedule never fires: + +| step | where | +|---|---| +| arm | `ci.yml`'s `plan` job enables it on a push to master that touches `hw/`, the configuration TOMLs, `VERSION`, the gate's scripts, catalog or baselines, or the workflow; an indeterminate diff arms it too | +| disarm | the scheduled run disables it as its first step, so a push that lands while it synthesizes re-arms the next night | +| retry | a build error re-enables it; a pass or a regression is a verdict, and its run is the record | -It pins master and skips an unchanged commit like `fpga_gate.yml`, but a -hosted runner keeps no state, so the marker is a cache entry keyed by the -SHA and the spec's hash. +A quiet night creates no run. A manual run needs the workflow enabled first +(`gh workflow enable asic_gate.yml`). The `toolchain` job exists because the cache key carries the commit the toolchain pin resolves to: refreshing the prebuilt toolchain misses it by