[Alex-Vasile]

JIRA ID : AIHPBLAS-5255

Motivation

The TensileLite common tests build GPU kernels for each YAML test config and then run them through tensilelite-client, the C++ program that runs and validates generated kernels. The changes below this one in the stack speed up both stages. This PR keeps the before and after timings of those changes and the scripts to measure a before and after again. It is meant to stay a draft, as a reference for reviewers of the other PRs.

Technical Details

All files are new and sit under projects/hipblaslt/tensilelite/timing_runs/. Nothing in the build, CI or tests reads them, and no existing file changes.

• Timing aggregates: 8 AGGREGATE_analysis.txt files, each the sum of one per-config pass (every config built or run on its own, with timers on), written by scripts/aggregate_timing.py. gfx950: full_runs_gfx950/baseline_build/ and baseline_run/ (before), after_run/ (after the run-stage changes) and after_build/ (after the build-stage changes). gfx942: full_runs_gfx942/pre_build/ and pre_run/ (before), post_build/ and post_run/ (after).
• Scripts, in full_runs_gfx950/, used on both GPUs: split/run_split.sh (the CI-like pass: whole build stage, then whole run stage, 4 pytest-xdist workers), per_config/run_per_config_build.sh and per_config/run_per_config.sh (the per-config passes), and measure_env.sh, which they all read for paths, caches, GPU target, cores and GPU. before_after/run_side.sh runs one side of a before/after comparison through all 3 passes, and before_after/compare_aggregates.py prints 2 aggregates side by side.
• run_split.sh takes DIST, the pytest-xdist scheduler (default worksteal, as the hw-common CI category uses), and CLEAR_RUN_CACHE (default 0: the run stage keeps the cache the build stage left). The CI-like numbers below were measured with DIST=load and CLEAR_RUN_CACHE=1.

The PR's first commit also added the exploratory runs, the investigation and decision notes and the write-ups; the second commit removes them, so they stay in the history.

How to read these results

Each result set was measured at one commit of the original, unsplit branch users/alvasile/new_test_client ("commit N" is its position among that branch's 109 commits, 8149ff02ff2d..bb98b8d3b0bb). The last column names the PRs of this stack that hold the same code changes.

Result set	Code measured (original-branch commit)	Equivalent point in this stack
gfx950 before: baseline_build/, baseline_run/ (184 configs)	Commit 35, 289f976d5a52; its code equals commit 20	Part of #15256 (the timers of commits 2-20) and all of #15262. None of the other PRs.
gfx950 after the run-stage changes: after_run/ (185 configs)	Commit 75, a9221990f377	#15255 to #15278, except commit 97 (the last commit of #15265) and commits 100, 103 and 105 of #15256. None of #15251-#15254 or #15279-#15290.
gfx950 after the build-stage changes: after_build/ (185 configs)	Commit 94, bcf018e67fd0; the first 25 configs at commit 92, 55f2c5569f5a. The client is commit 75's; no client source changed between 75 and 94.	The previous row plus #15252, #15253, #15281-#15285, #15286 without commit 106, and #15287-#15290
gfx942 before: pre_build/, pre_run/ (123 configs)	Commit 36, aa19af72bc03; its code equals commit 20	Same as gfx950 before
gfx942 after: post_build/, post_run/ (123 configs)	Commit 98, c1345541ac0c	#15252-#15290, except #15279, #15280, commit 106 of #15286, and commits 100, 103 and 105 of #15256

Reading notes:

• No PR of this stack was measured on its own: each before/after pair compares 2 commits of the original branch.
• The stack reorders that branch. The fixes (ci(tensilelite): run the hw-common tests with pytest-xdist worksteal #15251-fix(tensilelite): check the Synchronizer after every launch of a solution #15255) are commits 59-102 there, so no "before" point contains a fix. Most of the instrumentation (feat(tensilelite): time the client, the library-client step and the test artifacts, and report the timing by purpose #15256) and the reference change of perf(tensilelite): recompute the CPU reference only when its inputs were rewritten #15262 are commits 2-20 there, so every "before" point already contains them.
• gfx950 has a separate point after the run-stage changes; gfx942 does not, so its "after" holds both sets of changes. Each aggregate covers one stage, so the run aggregates give the run-stage gain and the build aggregates the build-stage gain.
• The gfx950 runs were made before the original branch was rebased onto 8149ff02ff2d, from develop af089c49d09c; the rebase changed DataInitialization.hpp, main.cpp and ReferenceValidator.cpp to resolve conflicts. This stack is based on a newer develop (e3f8cddfeb95) than any measured tree.

How to time a change

Build one msgpack-only client per side, from that side's source tree:

cmake --preset tensilelite -S <tree>/projects/hipblaslt -B <build> -G "Unix Makefiles" \
  -DCMAKE_BUILD_TYPE=Release -DGPU_TARGETS=gfx942 \
  -DCMAKE_CXX_COMPILER=/usr/bin/amdclang++ -DCMAKE_C_COMPILER=/opt/rocm/bin/amdclang \
  -DHIPBLASLT_ENABLE_YAML=OFF -DTENSILELITE_ENABLE_CLIENT=ON -DTENSILELITE_BUILD_TESTING=OFF \
  -DTENSILELITE_CLIENT_ENABLE_ROCPROFSDK=OFF
cmake --build <build> -j --target tensilelite-client

Then start both sides at the same time, each on its own 25 cores and GPU, so that both see the same host load:

S=projects/hipblaslt/tensilelite/timing_runs/full_runs_gfx950/before_after
ROOT=/scratch/cmp TL=<before tree>/projects/hipblaslt/tensilelite \
  CLIENT=<before build>/tensilelite/client/tensilelite-client \
  GPU_TARGET=gfx942 CORES=0-24 GPU=0 bash $S/run_side.sh before &
ROOT=/scratch/cmp TL=<after tree>/projects/hipblaslt/tensilelite \
  CLIENT=<after build>/tensilelite/client/tensilelite-client \
  GPU_TARGET=gfx942 CORES=48-72 GPU=4 bash $S/run_side.sh after &
wait
python $S/compare_aggregates.py /scratch/cmp/before/build /scratch/cmp/after/build
python $S/compare_aggregates.py /scratch/cmp/before/run /scratch/cmp/after/run

Each side's CI-like stage times are in /scratch/cmp/<side>/split/times.txt. The per-config passes cover the configs whose run passed in either side's CI-like pass. If the "before" tree is older than the timers you compare, set ANALYZE_DIR to the "after" tree's scripts/ directory so both sides are read by the same analyzer.

Test Plan

Run from projects/hipblaslt/tensilelite:

• pre-commit run --files on the added files checks whitespace, line endings and the hooks that apply to shell and Python.
• bash -n on the shell scripts checks their syntax.
• compare_aggregates.py on the gfx942 per-config data reproduces the comparison table removed from this PR.

Test Result

Time saved: none by this branch; it adds documentation, reports and scripts only. The headline numbers below measure the whole stack, not this PR.

All numbers are wall times of the CI-like pass (-n 4, timers off), measured on the original unsplit branch on shared hosts. Each run used 25 cores reserved with taskset, OMP_NUM_THREADS=25 and 1 GPU.

gfx942 (MI300X). Before and after ran at the same time on different sockets and GPUs, 1 sample each:

Measurement	Before	After	Change
Build stage	1,028 s	634 s	-38%
Run stage	1,494 s	742 s	-50%

gfx950 (MI355X). Before and after ran back to back in the same session, each on an idle GPU. For the build stage, "after" is the point after the build-stage changes; for the run stage, the point after the run-stage changes:

Measurement	Before	After	Change
Build stage	616.5 s	535.5 s	-13.1%
Run stage	2,445 s and 2,499 s	941 s and 943 s	-61.5% and -62.3%

On the added scripts: pre-commit passed, bash -n passed, and compare_aggregates.py reproduced the gfx942 build comparison row for row. The scripts were not run end to end on this branch. Automated GitHub checks are pending.

Related

• Work tracking: AIHPBLAS-5255

Device / Architecture Coverage

Documentation, generated reports and hand-run scripts only. No product code, test, CI configuration or GPU behavior changes, so passing PR checks are sufficient.

Submission Checklist

• Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

Risk level

Low (1/5): docs, data and hand-run scripts only; nothing builds, runs or is tested differently.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*). Current: code file(s) changed: projects/hipblaslt/tensilelite/timing_runs/full_runs_gfx950/before_after/compare_aggregates.py; no test file found
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[therock-pr-bot]

🎉 All checks passed! This PR is ready for review.