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/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/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/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/ci/testcases/core.yaml b/ci/testcases/core.yaml index 92532beef2..b09ebb949e 100644 --- a/ci/testcases/core.yaml +++ b/ci/testcases/core.yaml @@ -146,10 +146,6 @@ 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%) via: blackbox app: sgemv args: -m512 -n512 @@ -159,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 @@ -168,10 +163,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 2751838e92..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 - # 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%) via: blackbox app: dxa_copy_mcast configs: -DVX_CFG_EXT_DXA_ENABLE @@ -507,15 +505,12 @@ 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%)" 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 +518,48 @@ 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 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..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 @@ -187,7 +182,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 +213,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..4cc78dbf65 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,24 +207,28 @@ 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 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 - 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 +239,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..04238ea748 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` @@ -370,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 | @@ -422,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` @@ -443,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 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/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/alu_unit.cpp b/sim/simx/alu_unit.cpp index 648729c0ad..a866e2f1ca 100644 --- a/sim/simx/alu_unit.cpp +++ b/sim/simx/alu_unit.cpp @@ -25,9 +25,16 @@ 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) {} +// 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 +52,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,25 +69,25 @@ 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) { + // 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: 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; + return kMulDivLatency; default: std::abort(); } @@ -562,8 +569,13 @@ void AluUnit::on_tick() { if (!output.full()) { auto trace = input.peek(); this->execute(trace); - uint32_t delay = this->latency_of(trace); - output.send(trace, delay); + // 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; + } + 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 b5f5288fba..6e2db38121 100644 --- a/sim/simx/alu_unit.h +++ b/sim/simx/alu_unit.h @@ -19,8 +19,18 @@ 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. + SimChannel branch_ctl_out; + protected: void on_tick() override; diff --git a/sim/simx/core.cpp b/sim/simx/core.cpp index 61e80e974c..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]; @@ -122,18 +125,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 @@ -167,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). @@ -178,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) @@ -192,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 @@ -231,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_); @@ -249,19 +245,25 @@ 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_); 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. { @@ -272,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 < @@ -280,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)); } } @@ -294,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; @@ -315,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(); @@ -518,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(); @@ -530,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); + } } } @@ -539,7 +580,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); @@ -551,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)) { @@ -568,12 +618,21 @@ 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. + 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; + } 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 + // 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 // A TRACE macro must hold a ray-pool slot before its head uop enters @@ -586,26 +645,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(); @@ -620,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); @@ -676,53 +727,59 @@ 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 + // 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 // 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) { @@ -732,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 @@ -756,16 +814,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); @@ -785,6 +854,9 @@ class Core::Impl { scoreboard_->release(trace); } } + if (trace->eop && trace->instr_ptr->fcsr_writes()) { + scoreboard_->release_fcsr(trace); + } if (trace->eop) { @@ -859,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(); @@ -1032,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_; @@ -1063,6 +1143,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/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 30f6109f91..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) {} @@ -41,6 +44,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) { @@ -51,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; @@ -105,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); @@ -118,16 +132,18 @@ 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 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/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/dxa/dxa_core.cpp b/sim/simx/dxa/dxa_core.cpp index 3f51c3254e..5f3c357644 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); @@ -635,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 @@ -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; @@ -709,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; @@ -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/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/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/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 bf15096734..63a275bba8 100644 --- a/sim/simx/instr_trace.h +++ b/sim/simx/instr_trace.h @@ -67,9 +67,12 @@ 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 ; + // Cycle the instruction entered its warp's instruction buffer. + uint64_t ibuf_time; instr_trace_t(uint64_t uuid) : uuid(uuid) @@ -93,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) {} @@ -120,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/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; 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/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/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..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,11 +35,34 @@ 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(); +} + +void Scoreboard::on_tick() { + uint64_t now = SimPlatform::instance().cycles(); + while (!pending_releases_.empty() && pending_releases_.front().due <= now) { + auto& r = pending_releases_.front(); + 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()) { + this->tick_sleep(); + } } 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)) { @@ -89,9 +114,24 @@ 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, 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(); } bool Scoreboard::commit_packet(instr_trace_t* trace) { diff --git a/sim/simx/scoreboard.h b/sim/simx/scoreboard.h index 4eeece71fc..e05ade6479 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 @@ -50,15 +54,33 @@ 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(); 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 = 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/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/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..5e2bd89173 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,17 @@ class TcuUnit::Impl { } #endif exec_done_.fill(false); - wgmma_planned_warps_.fill(0); + for (auto& due : result_due_) { + due.clear(); + } 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,131 +559,34 @@ 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 + 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()); - #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 +615,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 +635,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 @@ -755,19 +696,8 @@ class TcuUnit::Impl { } #endif if (simobject_->Outputs.at(b).try_send(trace, delay)) { + due.push_back(now + 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 +708,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 +1384,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 +1462,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 +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_; - // Per-block bitmask of warp IDs with planned WGMMA lines; cleared on fu_unlock. - std::array wgmma_planned_warps_; + // 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(). @@ -1366,9 +1575,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; diff --git a/sim/simx/types.h b/sim/simx/types.h index a0133d53e1..6423d1a61a 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)); } @@ -1971,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_; 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.