[Alex-Vasile]

JIRA ID : AIHPBLAS-5256

Motivation

When TensileLite benchmarks a config (the path every common test and every tuning run takes), each benchmark step generates assembly kernels in a process pool, then assembles them, then compiles Kernels.cpp, the C++ source of the step's helper kernels, with clang. That helper compile was 1,004.7 s, 39% of a gfx950 per-config build aggregate, and it started only after codegen and assembly, although its input is ready at the start. Each step also filled and drained its own pool, so configs of many small steps left most workers idle. A third cost: the invalid-kernel check ran in its own pool, which pickled every kernel's state to a worker and back for a one-attribute test.

Technical Details

TensileCreateLibrary/Run.py replaces writeSolutionsAndKernels with 3 phases:

• prepareSolutionsAndKernels: setup, write Kernels.cpp/Kernels.h, and queue the helper compile on helperCompileExecutor(), a 1-thread executor that compiles steps one at a time, in step order, so a later step still hits the helper cache.
• generateBenchmarkKernels: 1 ParallelMap2 pool over the kernels of any number of steps. Sources held across steps are compressed with the existing memcompress.
• finishSolutionsAndKernels: invalid-kernel check, now a list comprehension in the parent, then assemble, link, and wait for the helper compile (timer python_kernel_build_src_co_wait).

benchmarkPoolSize() gives every pool of the run the CPU budget minus 1, reserving 1 slot for the compile; with a budget of 1 the compile finishes before codegen. One size per run avoids loky, joblib's process-pool backend, restarting workers when the pool size changes.

BenchmarkProblems.py: main prepares every step of a config, generates all their kernels in 1 pool, then finishes each step in config order as before (validate, assemble, library, client run, results). A step waits for the batch to finish first only when it reads what earlier steps write: the same cache directory, or a repeated problem-size group name. The Ductile backend, a search that reads each evaluation's results to pick the next candidates, keeps stepping one at a time through a new OptimizationBackend.uses_benchmark_results() (default True; TensileBackend returns False).

scripts/analyze_timing.py puts the wait in the timer tree and moves python_kernel_build_src_co to the overlapping totals so it is not double counted.

New tests: test_benchmark_helper_compile_overlap.py and test_benchmark_steps_one_pool.py (17 cases), plus TestHelperCompileOverlap in test_analyze_timing.py; 4 existing test files follow the new names.

Test Plan

Run from projects/hipblaslt/tensilelite:

• tox -e unit -- tensilelite/Tests/unit/test_benchmark_steps_one_pool.py tensilelite/Tests/unit/test_benchmark_helper_compile_overlap.py tensilelite/Tests/unit/test_benchmarkProblems_solution_pool.py tensilelite/Tests/unit/Ductile tensilelite/Tests/unit/test_codegen_worker_state.py tensilelite/Tests/unit/test_analyze_timing.py tensilelite/Tests/unit/characterization/TensileCreateLibraryRun checks the phase order, overlap, wait rules, Ductile stepping, error handling and the timing report.
• tox -e py3 -- tensilelite/Tests -m common -k 'custom_mainloop_scheduling or spmm_i8_sb or dtl' on a gfx950 GPU builds and runs configs with many steps.
• pre-commit run --files on the changed files checks formatting and lint.

Test Result

Time saved: overlapping the helper compile cut build wall time by 12% on fp16_tn, 6% on subtile_mxfp4 and 2% on dtl; one pool per config then cut it by a further 58% on custom_mainloop_scheduling (72.3 to 30.0 s), 32% on subtile_mxfp4, 23% on spmm_i8_sb, 18% on f8_gate_r and 8% on dtl. Each was measured on the original unsplit branch against the commit before it, --build-only on the gfx950 (MI355X) host, 24 cores reserved with taskset (budget 24), cold helper cache, 2 runs per side. Moving the invalid-kernel check to the parent saved no wall time: CI-like gfx942 (MI300X) build stage, 25 cores, -n 4 --dist=worksteal, 441 s against 441.5 s, mean of 2 runs per side.

The individual runs:

config (steps)	build before	overlap only	overlap and one pool
fp16_tn (3)	50.5, 50.6 s	44.4, 44.7 s	not measured
custom_mainloop_scheduling (66)	not measured	71.5, 73.1 s	30.0, 30.0 s
spmm_i8_sb (8)	not measured	29.8, 29.8 s	23.3, 22.7 s
dtl (26)	142.8, 142.0 s	139.1, 139.1 s	129.8, 127.0 s
subtile_mxfp4 (130)	83.2, 83.0 s	78.1, 78.4 s	53.0, 52.9 s

The table joins 2 separate before/after measurements, one per source commit. Where both measured the middle column, it shows the second; the first gave dtl 139.5, 139.6 s and subtile_mxfp4 77.7, 78.5 s.

• Artifacts of these configs matched the parent's member by member, apart from clang's random __hip_cuid_* symbol in the helper code objects. Solution state after each step's build and every output file matched for 3 configs.
• Pool processes plus the helper compile stayed at or below 24 in every run. A compiler launcher forced to fail made the test fail with the compiler error, finished the earlier steps and left no process behind.
• Peak memory of the Tensile process: dtl 2.8 to 1.5 GB. All processes together: custom_mainloop_scheduling 6.9 to 8.5 GB, because all workers generate at once.
• Moving the invalid-kernel check to the parent: python_kernel_validate 124.1 to 27.7 s summed over the 4 workers, while assembly time rose by 71.4 s. The unit suite there had the same 519 failures as its parent, in test_custom_kernel_staggeru_metadata.py and test_subtile_ds_read_emit.py.

Not yet run on this branch. Automated GitHub checks are pending.

Related

• Work tracking: AIHPBLAS-5256
• #13671 overlapped the helper compile with assembly codegen on the TensileCreateLibrary library path; this PR does the same on the benchmark path.

Device / Architecture Coverage

This changes the benchmark kernel-build pipeline: order, concurrency and pool sizes; generated kernels and helper sources are unchanged. The overlay requires CI coverage of kernel-generation changes on the affected archs, gfx942 and gfx950 at least. An arch that is only built counts as uncovered.

Submission Checklist

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

Risk level

High (4/5): restructures the benchmark build used by every common test and tuning run. The order of work and log output changes. A helper-compile failure is raised when its step finishes, and an error in the shared pool now ends the run before earlier steps are built. The Tensile process holds every step's compressed sources until they finish. Coverage so far is from the original branch only.

[therock-pr-bot]

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 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.