[tony-davis]

JIRA ID : AIHPBLAS-5352

Motivation

The gfx1250 hardware workflow runs the hipBLASLt client tests on a host shared by 4 runners. Since #14545, the 93 large-memory cases (2^31- and 2^32-element sizes) run after the 4 test shards, one group per process, and behind a lock file shared by every runner on the host (flock on /ci-locks), so only one job at a time runs them. That lock was a stopgap: without it, several jobs' large cases together ran the host out of memory, because the test client kept their pinned host buffers until the process exited. The three PRs below this one fix that in the client, and #14569 gives each job's container its own 110 GiB memory cap, so the lock can go.

I don't think this is a guarantee that the host never runs short, and I'm not planning a long test campaign to prove it. The client now holds much less pinned memory and skips a case with the amounts instead of getting killed, and the caps keep any kill inside one job. That's better than a lock that serializes every job's large phase. If kills come back, putting the lock back is a revert of this one file.

Technical Details

• The workflow no longer takes the lock, and the /scratch/ci-locks mount is removed. Jobs on the host run their large phases at the same time.
• The large cases still run after the shards, one group per process, and the shards still exclude them. Each shard checks memory only for itself and at one moment, so several shards can pass the check together; until that is measured (fix 6 in AIHPBLAS-5352), the split stays.
• The large phase sets HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1. These cases are the reason the job runs on this host, so one that skips for lack of memory now fails the job instead of passing quietly. The shards keep the default, where a memory skip is a skip.
• hipblaslt-test now prints its own peak resident memory at the end of a run, as [ PEAK RSS ] N.NN GiB of host memory (Linux). The job summary reads that line into a table of every process's peak, which counts pinned host buffers, and the number of its tests that stopped for memory, followed by the [ MEMORY ] lines themselves. A process that dies before the end shows no peak.
• The peak comes from the client rather than GNU time, because the estimate check (HIPBLASLT_TEST_CHECK_MEMORY_ESTIMATE, from fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610) resets the kernel's peak, which also resets what time and getrusage report: a process that touched 1 GiB and then reset it reported 2 MiB. MemoryEstimateCheck (memory_estimate_check.hpp) now saves the peak before each reset, and hipblaslt_gtest_main.cpp prints the larger of that and the final peak. This also drops time from the job's packages and its unscrubbed timing files.
• TESTING.md's row for this lane describes the large phase and the summary.

This is the fourth PR in the AIHPBLAS-5352 stack and depends on #14610, #14609 and #14608. It changes the workflow, TESTING.md, and for the peak, memory_estimate_check.hpp, hipblaslt_gtest_main.cpp and matmul_gtest.cpp. It relies on the client behavior in the PRs below it, so it merges after them.

Test Plan

• actionlint on the workflow: the same 3 findings as develop and none new (a custom runner label, and two deliberate glob matches in the existing in_filter).
• pre-commit run on both files.
• The edited step's large-phase and summary script, extracted from the workflow and run locally against a stub hipblaslt-test: once with every case passing, and once where a large case skips for memory. The summary block alone under bash -e, against one log with a [ PEAK RSS ] line and one without.
• MatmulMemory_pre_checkin.peak_resident_memory_survives_the_estimate_check pins 1 GiB, frees it, runs the estimate check, and expects the kernel's peak to drop while the reported one keeps the 1 GiB. A variant that does not save the peak fails it.
• On gfx1250, this PR's own runs of the hardware workflow, which test the whole stack without the lock. After that, the lane's normal runs are the test.

Test Result

• Stub run, every case passing: exit 0, and the summary table lists the 4 shards and 6 large groups. A memory skip in a shard shows up in the table and the list without failing the job.
• Stub run, a large case short of memory: exit 1, and that case is listed in the summary.
• Summary block: a log with [ PEAK RSS ] 12.34 GiB shows 12.34; a log without the line gets an empty cell, and the step goes on to its exit code.
• The peak test passed 5 of 5 runs. On gfx90a the memory suites print [ PEAK RSS ] 3.19 GiB of host memory.
• gfx1250, first run without the lock (41717824185, Archana's read): red only from the already tracked intermittent AmaxD case. No lock wait, 0 kills, no memory skips; [ PEAK RSS ] 15.3 to 38.8 GiB per shard and 2.2 to 49.9 GiB per large group; TensileLite 216 passed. Its large phase overlapped about 4 minutes with another job's, so 2 at once, not 4. Before this PR, her dispatch of the stack with the lock (41582794047) passed hipBLASLt twice with no kills and no memory skips.
• Automated GitHub checks are pending.

Submission Checklist

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

Risk level

Medium (3/5). It removes the only coordination between jobs on the gfx1250 host. What bounds memory now: each container's 110 GiB cap from #14569, which keeps an out-of-memory kill inside one job, and the client's preflight and buffer release from this stack. Four containers at 110 GiB plus about 36 GiB the host uses when idle come to 476 of about 499 GiB, about 23 GiB to spare (Archana's numbers). The caps hold for pinned buffers too: they are shared memory charged to the container. A large case that runs short of memory now fails the job, so a problem shows up as a red run, not a quiet skip. The lane is not a required check, and putting the lock back is a revert of this file. The hipBLASLt library is unchanged; the test client only gains the peak report.

Related

• Work tracking: AIHPBLAS-5352 (acceptance: drop the host-wide lock once the cap and the harness changes are in, and keep the cap).
• Depends on fix(hipblaslt): run the memory preflight for every matmul case and list memory skips #14610, fix(hipblaslt): make the test client's host-memory query respect container limits #14609 and fix(hipblaslt): keep large pinned host buffers only while tests reuse them #14608. Replaces the lock from ci(gfx1250): run the large-memory hipBLASLt tests one job at a time #14545; relies on the cap from feat(ci): cap the gfx1250 test container's memory #14569.
• Follow-up on the harness design: AIHPBLAS-5451.

Change class

build/ci, ci-gate-change (a memory skip in the large phase now fails this lane). No new workflow, runner, trigger or shared CI state; the checkout is unchanged.

[Copilot]

🟡 Changes recommended

The workflow publishes unsanitized timing data, can under-report peak memory after measurement resets, and aborts when the timing file is missing even if all tests pass.

3 open issues

• Timing logs expose measurements that the workflow should scrub · New
• Peak memory can be under-reported after clear_refs resets · New
• Missing time utility makes the summary assignment abort the step · New

What changed in this PR

Removes host-wide serialization of gfx1250 hipBLASLt large-memory tests, relying on the container memory cap and client memory-management fixes from the prerequisite PRs.

Changes:

• Removes the shared lock and mount while retaining the four-shard phase followed by separate large-test groups.
• Makes memory skips fail the large-test phase and adds per-process memory reporting.
• Updates the lane documentation to describe execution order, failure behavior, and summaries.

File	Description
projects/​hipblaslt/​TESTING.md	Documents large-test execution, memory-skip failures, and reporting.
.github/​workflows/​gfx1250_hipblaslt_tensilelite_test.yml	Removes host locking and adds strict memory-skip handling and memory summaries.

_{Review effort: Balanced}

[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?

[codecov]

Codecov Report

✅ All modified and coverable lines are covered by tests.

Additional details and impacted files

@@                                 Coverage Diff                                  @@
##           users/tony-davis/AIHPBLAS-5352-memory-preflight   #14982       +/-   ##
====================================================================================
- Coverage                                            77.65%   70.05%    -7.59%     
====================================================================================
  Files                                                 2905       67     -2838     
  Lines                                               490744    22641   -468103     
  Branches                                             74315     3541    -70774     
====================================================================================
- Hits                                                381040    15861   -365179     
+ Misses                                               84683     5279    -79404     
+ Partials                                             25021     1501    -23520     

Flag	Coverage Δ
TensileLite-CPP	?
TensileLite-Unit	?
hipBLAS	?
hipBLASLt	70.05% <ø> (ø)
hipCUB	?
hipDNN	?
hipFFT	?
hipRAND	?
hipSOLVER	?
hipSPARSE	?
rocBLAS	?
rocFFT	?
rocRAND	?
rocSOLVER	?
rocSPARSE	?
rocThrust	?
see 2838 files with indirect coverage changes

🚀 New features to boost your workflow:

• ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.
• 📦 JS Bundle Analysis: Save yourself from yourself by tracking and limiting bundle sizes in JS merges.

[archana-ramalingam]

Thanks @tony-davis! Both #14610 asks are in (HIPBLASLT_TEST_FAIL_ON_MEMORY_SKIP=1 on the large phase, [ PEAK RSS ]). The large phase still runs after the shards, so dropping the lock can't raise a container's peak; what's left is host-wide.

First lock-free run 41717824185: red from the already tracked intermittent AmaxD case.
No lock wait, 0 kills, no [ MEMORY ]; [ PEAK RSS ] 15.3–38.8 GiB per shard, 2.2–49.9 per large group; TensileLite 216 passed / 0 failed.
All 4 runners were busy, but the large phase is ~8 of a ~50-min job: ours overlapped ~4 min with #14123's (41730286382, with the lock) while the other two GPUs were in TensileLite. So 2 at once, not yet 4.

Sources, all on the gfx1250 host: free -g and df -h /dev/shm (11-03); the job image under docker run --user 0:0 --memory 3g (11-04); our 10-28 probe of develop 2a5d491 before the stack (4 shards under /usr/bin/time -v, host memory sampled every 10 s).

Non-blocking:

1. Risk line: 4 × 110 GiB + ~36 GiB idle = 476 of 499 GiB, ~23 GiB margin. The caps hold for pinned buffers too: they're anonymous shared memory charged to the container (probe: 333.6 GiB shmem, more than /dev/shm holds; 8.6 MB after each run), and /dev/shm has 1.9 MB after the 11-02 kills.

2. Test plan, if you want the 4-overlap run: same-branch dispatches share a concurrency group (event_name-ref, cancel-in-progress, #L114-L115) and cancel each other (this PR's 41706411592/41706411814/41717823133), and staggered runs rarely line up anyway. To force it: push the head to 4 refs and, with the runners idle, dispatch all 4 with run_tensilelite=false and hipblaslt_gtest_filter='*fast_check_size_t_index_overflow*:*stress_matmul_size_t_index_overflow*:*size_t_output_overflow*:*offset_above_2_32*:*batch_offset_large*'. The shards run nothing, so the 4 large phases start after their 3–5 min builds and overlap for most of their ~8 min.

3. To measure the shard race (fix 6 in AIHPBLAS-5352): the job container mounts /sys/fs/cgroup read-only, so memory.peak can't be reset and at job end includes the build's page cache. Sample the non-reclaimable part instead:

   ( while sleep 5; do awk '/^(anon|shmem) /{s+=$2} END{print s}' /sys/fs/cgroup/memory.stat; done > "$OUT/cgroup_mem.txt" ) & SAMPLER=$!
   # ... shards and large phase ...
   kill $SAMPLER; echo "container anon+shmem peak: $(sort -n "$OUT/cgroup_mem.txt" | tail -1)"; grep -E '^(max|oom_kill) ' /sys/fs/cgroup/memory.events
   

   Per-process peaks can't show it: 134 GiB summed vs ~69 GiB together.

4. Update PR description:

   • since 9f7bcf1 the peak comes from [ PEAK RSS ], not /usr/bin/time -v (Technical Details, Test Result)
   • PR also changes memory_estimate_check.hpp, hipblaslt_gtest_main.cpp and matmul_gtest.cpp.
   • Update hardware run link: 41717824185

[tony-davis]

Thanks Archana, for the approval and for all the digging! Your gfx1250 numbers have been a huge help.

I updated the description with your points: [ PEAK RSS ] instead of time -v, the client files this PR touches, your lock-free run, and your margin math in the risk line.

On the 4-way overlap run and the cgroup sampler, I think I'm going to stop here, and I hope that's OK. My view is that this stack doesn't guarantee the host never runs short. It just makes things a lot better: the client holds much less pinned memory, gives it back sooner, and skips with the amounts instead of getting killed. That's the part I can fix from the client side. If a big test campaign showed it wasn't enough, I don't think I'd do anything different here. I'd put the lock back (a one-file revert) and pick up fix 6, and the caps keep any kill inside one job in the meantime. So I'd rather let the lane's normal runs tell us.

Your sampler recipe is great, though, so I'll put it on AIHPBLAS-5352 under fix 6 for whoever picks that up.

[newling]

Small PR with the correct shape, title and description look straightforward. Approving (but haven't been through in detail, I think Archana has)

[therock-pr-bot]

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