Repository navigation
test(hipblaslt): run applicable client tests on gfx1250 and gfx1250-strict - #13504
Merged
Merged
Conversation
…rray values gpu_arch and gpu_arch_exclude were 16-byte fields, which can't hold a pattern such as '9(0a|42|50)|1250'. Widen both to 32 bytes in the Arguments struct and in hipblaslt_common.yaml. hipblaslt_gentest.py also stored a char-array value that exactly filled its field with no terminator, so the C++ side read past it. Generation now fails with the field name and limit for any value that doesn't fit with its NUL. Adds CPU-only unit tests for that check. Co-authored-by: Cursor <cursoragent@cursor.com>
…trict Review every architecture-filtered client test case for gfx1250. Enable the 33 cases whose problem types gfx1250 ships kernels for (fast_check and other address-arithmetic cases, xf32, OCP FP8 scales and mixes, i8, f64, bf16 bias gradient, multithread streams, GSU tuning), and fix four grouped-GEMM cases whose gfx11/gfx12 exclusion also excluded gfx1250. Cases kept off gfx1250 because it has no kernels for their problem type carry a "Not on gfx1250" comment. Document how gpu_arch, gpu_arch_exclude and known_bug_platforms treat gfx1250 and gfx1250-strict in TESTING.md. Co-authored-by: Cursor <cursoragent@cursor.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #13504 +/- ##
===========================================
- Coverage 75.95% 73.11% -2.85%
===========================================
Files 2880 2865 -15
Lines 479080 474107 -4973
Branches 71971 70817 -1154
===========================================
- Hits 363883 346614 -17269
- Misses 90339 103169 +12830
+ Partials 24858 24324 -534
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
tony-davis
marked this pull request as ready for review
October 16, 2026 20:36
The gfx1250-strict hardware run on this PR failed this case with "NO solution found!" from the all-solutions find, so gfx1250 has no kernel for this problem. Restore the develop arch pattern (1201) and note why. Co-authored-by: Cursor <cursoragent@cursor.com>
archana-ramalingam
approved these changes
Oct 21, 2026
archana-ramalingam
left a comment
archana-ramalingam
left a comment
Contributor
There was a problem hiding this comment.
Please update PR description to reflect latest gfx1250 HW CI results.
Reviewed minor nits below.
Resolve the TESTING.md conflict by keeping develop's gfx103X/gfx120X row and this PR's gfx1250 row. Update the gfx1250 row and the scoping notes to match the current gfx1250 HW workflow (per-PR, nightly, manual dispatch), the UTF-8 byte limit on gpu_arch fields, and the unanchored `|1250` pattern, per review feedback. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
Thanks James! I took the TESTING.md wording from your docs commit. I'm deferring the generator CLI tests and the pre-commit hook to a follow-up so this PR stays test-data and harness only. |
Contributor
Author
|
I dispatched a manual run of It uses |
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
tony-davis
pushed a commit
that referenced
this pull request
Oct 28, 2026
…14545) JIRA ID : N/A ## Motivation The gfx1250 HW host ran out of memory (10-23, 10-27); the workflow is disabled. Cause: #13504 enabled 93 hipBLASLt tests with 2^31/2^32 sizes on gfx1250. All 4 shards reach them at once, and the client's pool keeps their pinned host buffers, so one job nearly fills the host. ## Technical Details - Shards skip the 93 (5 name patterns). - After the shards, the 93 run in 6 processes, one per group (the two stress groups apart; largest ~60 GiB), behind a host-wide `flock` (`/scratch/ci-locks`), one job at a time; 50-min wait cap. - A custom `hipblaslt_gtest_filter` gets the same split: the shards add the 93 to its negatives, and the large step runs the large tests the filter selects (from `--gtest_list_tests`). - Fails if a large group matches no test (renamed test) or a large log runs 0 tests. - Stopgap until the harness bounds its own host memory (AIHPBLAS-5352); then remove the split and the lock. - Same tests, same report: the large logs are `hipblaslt_test_shard_large*.log`. ## Test Plan - Step script with a stub `hipblaslt-test`. - On the CI host, runners stopped: the patterns match exactly the 93; the full suite; 4 concurrent jobs with the lock. - This PR's HW run once the workflow is re-enabled. ## Test Result - Single job: 69,569 + 93 passed, 0 failed. - 4 concurrent jobs: the lock serialized the large phases; host memory peaked under half, no OOM kills. One unrelated wrong-result failure on one GPU, tracked separately. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. gfx1250 HW workflow only; same tests, different order. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude <noreply@anthropic.com>
1 task done
tony-davis
added a commit
that referenced
this pull request
Nov 6, 2026
… them (#14608) JIRA ID : AIHPBLAS-5352 ## Motivation `hipblaslt-test`, the hipBLASLt test client, keeps the host buffers that each test frees in a pool, so the next test can reuse them instead of allocating again. These buffers are pinned (allocated with `hipHostMalloc`, so the operating system can neither page them out nor reclaim them), and the pool kept every one until the process exited. On the gfx1250 hardware test host, which has about 499 GiB of memory, the 93 cases with 2^31- and 2^32-element sizes enabled in #13504 left about 123 GiB pinned in one process. With 4 CI runners each splitting the suite into 4 parallel test processes, the out-of-memory killer took the runners down on 10-23 and 10-27. Separately, a failed host allocation returns a null buffer, and the test then writes through it. I don't think this stack guarantees the gfx1250 host never runs short. What it does is make the client hold much less pinned memory, give it back sooner, and skip a case with the amounts instead of getting killed, which is the part of the problem the client can fix. If pressure comes back, the 110 GiB container caps from #14569 keep a kill inside one job, and putting the lock back (#14982 removes it) is a one-file revert. So I'm not planning a long soak test before merging. The pool comes from ROCm/hipBLASLt#1202, which reused allocations to speed up `hipblaslt-bench` across about 500,000 small matmuls. Reuse still matters for large buffers: on the test host below, `hipHostMalloc` takes about 0.18 s per GiB and `hipHostFree` about 0.06 s per GiB. ## Technical Details - A pooled host buffer of 1 GiB or more (`kHostPoolRetainLimit` in `d_vector.hpp`) now stays cached only while consecutive tests reuse it. The pool records the test count when a buffer comes back, and `HostPoolReleaseListener`, a new hook that Google Test calls after every test, frees the large buffers that the test just ended did not use. A run of same-shape cases, such as the four transpose combinations (NN, NT, TN, TT) of one size, still shares one allocation. - When a test asks for a large buffer that no cached buffer fits, the pool first frees the large buffers that earlier tests left, then pins new memory. This keeps the peak near one test's footprint when the shape changes. - A host buffer no longer gets the pool's 20% extra capacity when the extra would take it to 1 GiB or more (`host_pool_capacity()`), so a buffer under the limit stays under it. The existing cutoff for skipping the extra is a quarter of device memory, so on gfx1250 nearly every host buffer got it. - `HipHostBuffer::memcheck()` reports a failed pinned allocation, and `testing_matmul` checks every host buffer after allocating, including the four temporary buffers a swizzled A or B is built in. A failure now skips the test with the requested size ("could not allocate N GiB of pinned host memory"). - A pooled host buffer under 1 GiB is released once 100 consecutive tests have not reused it (`kHostPoolSmallIdleTests`). The pool reuses a buffer only for requests between half its size and its size, so without this a run of shrinking requests kept every earlier buffer pinned. - In `hipblaslt-bench`, a failed host allocation now exits with a failure, as a failed device allocation already did, instead of reporting success. - Device memory behaves as before. `hipblaslt-bench` does not install the listener, so it never frees large buffers early; it does see the 20% change for host buffers of 1 GiB or more. The hipBLASLt library itself is unchanged. This is the first of four stacked PRs for AIHPBLAS-5352. #14609 makes the client's host-memory query respect container memory limits, #14610 runs the memory preflight for every matmul case, and #14982 drops the host-wide lock in the gfx1250 workflow. ## Test Plan - `HostMemory_pre_checkin` (new, `clients/tests/src/host_memory_gtest.cpp`, 6 tests; the `_pre_checkin` suffix puts it in the standard tier the PR checks run). The tests check that a large buffer survives one test boundary and is freed after the next, that a same-shape test reuses it (what the first test wrote is still there), that a large miss frees two buffers earlier tests left, that a 1 GiB request gets exactly 1 GiB, that a small buffer is released after 100 tests that do not reuse it (and kept when one does), and that a failed allocation leaves a null buffer whose size is reported. A 0.9 GiB request gets exactly 0.9 GiB (the old rule gave it 1.08 GiB, over the limit). The tests skip, naming the buffer, on a host that cannot pin that much, and the reuse test skips on a Windows host where the pool frees its idle buffers before every request. A temporary variant without the small-buffer aging failed the new test. - To confirm the tests can fail, I built two temporary variants of the pool. One undid four behaviors at once (no release at the end of a test, no release on a miss, the 20% extra restored, `memcheck()` always succeeding), and it failed 4 of the 5 tests. The other freed large buffers whether or not a test used them, and it failed 3. Every test failed under at least one variant. - The 82 large-memory cases that apply to gfx90a (the five name patterns from #14545), and the `smoke` and `quick` tiers, on develop and on this branch, for wall time and peak resident memory. - Handle leaks: #14714 and #14700 on develop give the matrix layouts, matmul descriptors and HIP events owners, so an early return no longer leaks them. To check that with this stack's new early returns, a preloaded library counted creates and destroys of all three, on the full stack, over 1,892 matmul cases, once normally and once under a 1 MiB host limit so that most cases skip at the preflight. - Adjacent: the `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin` suites, including `memory_shortfall_reclaims_idle_buffers`, which uses the pool's existing `ReleaseCached()`. `hipblaslt-bench` was built, because it compiles the non-gtest form of the new check, but not run. ## Test Result One MI210 (gfx90a) on a shared host with 2.2 TB of memory. This branch's clients were built against the 10-25 TheRock nightly hipBLASLt gfx90a device library. Each pair of runs went back to back, but the host was busy (load average near 300), so the times carry some noise. | Measurement | develop | this PR | | --- | --- | --- | | 82 large-memory cases, wall time | 316.9 s | 313.8 s | | 82 large-memory cases, peak resident memory | 81.9 GiB | 49.2 GiB | | 82 large-memory cases, mean resident memory | 66.9 GiB | 16.6 GiB | | Resident memory after the last large case | 74 GiB | 2 GiB | | `smoke` + `quick` (9,171 tests), two runs | 66.6 s, 62.0 s | 63.6 s, 64.1 s | - Large-memory cases: 82 of 82 passed on both. `smoke` + `quick`: 9,169 passed and 2 skipped (`GroupedMatmul.smoke_Fp16NN` and `smoke_Bf16NN`, which skip on this GPU) on both. - `HostMemory_pre_checkin`, `FastCheck_pre_checkin` and `FastCheckDevice_pre_checkin`: 64 of 64 passed. - A first version that freed every large buffer at the end of each test took 449.6 s for the large cases against 318.2 s on develop, which is why the pool keeps a buffer while the next test reuses it. - Handle leaks: 0 layouts, descriptors or events left, both normally (13,580 created) and with 409 cases skipped at the preflight (11,235 created). - With the small-buffer aging, `smoke` + `quick` passed 9,225 tests in 59.3 s with peak resident memory 2.18 GiB. - gfx1250: this PR alone passed the hardware workflow's hipBLASLt tests in [40934683157](https://github.com/ROCm/rocm-libraries/actions/runs/40934683157), and Archana dispatched the full stack in [41582794047](https://github.com/ROCm/rocm-libraries/actions/runs/41582794047), which passed hipBLASLt twice with no kills and no memory skips. Her run on the rebased stack, [41717825380](https://github.com/ROCm/rocm-libraries/actions/runs/41717825380), was clean too. All three kept the host-wide lock; #14982 covers the run without it. - Automated GitHub checks are pending. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Medium (3/5). Every `hipblaslt-test` case passes through the new listener and the host pool, so a mistake here would affect the whole suite. The hipBLASLt library and device behavior are unchanged; only the test and bench clients change. The bench sees the 20% change near 1 GiB, and now exits with a failure when a host allocation fails. ## Related - Work tracking: [AIHPBLAS-5352](https://amd-hub.atlassian.net/browse/AIHPBLAS-5352), fixes 1 (pool retention) and 4 (null host buffer) of its list. - Cause: #13504. CI workaround this work is meant to replace: #14545, merged on 10-28. - Where the pool came from: ROCm/hipBLASLt#1202. Windows already clears the pool on smaller-memory hosts for the same reason: #4793, #6121. ## Device / Architecture Coverage Host-side test-client code only, with no kernel, library or kernel-selection change, so it does not depend on the GPU architecture and the standard PR checks are enough. Measured on gfx90a. ## Flags / Guardrails None. The 1 GiB threshold is the constant `kHostPoolRetainLimit`. [AIHPBLAS-5352]: https://amd-hub.atlassian.net/browse/AIHPBLAS-5352?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ --------- Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : AIHPBLAS-5134
Summary
About two thirds of hipBLASLt's architecture-filtered client test cases never ran on gfx1250: 132 of 196 entries in
matmul_gtest.yaml,smoke_gtest.yamlandrocroller_gtest.yaml. Most were scoped to gfx942 or gfx950, including the fast_check address-arithmetic cases. This PR reviews each one against what the gfx1250 libraries actually provide, and enables those that apply.1250added to theirgpu_archpattern, and 4 grouped-GEMM cases get a corrected exclusion. Each was checked type by type and layout by layout against the gfx1250 logic files (library/.../Logic/asm_full/gfx1250/), so every enabled combination has kernels there.size_t_index_overflow,size_t_output_overflow_epilogue,offset_above_2_32,batch_offset_above_2_32,store_extent_ROCM31016,size_t_output_overflow_all_solutions,every_solution_repeat_xf32,large_k_fp8,mixed_fp8, the five FP8 scale cases,matches_unit_check_bias_grad_bf16,threshold_custom_kernel_ROCM31258.size_t_index_overflow{,_hpa},size_t_output_overflow_bf16_ROCM26455,size_t_batch_stride_overflow,integer_exact_custom_kernel_store_extent_xf32_ROCM31258,batch_offset_large.gemm_double(and its smoke case),gemm_i8_dst_i32,gemm_i8_dst_i8,gemm_i8_dst_i32_alpha_float_bits,ext_i8_dst_i32_support_smoke,bgradb_bf16_NT,stride_lt_ld,multithread_stream,extapi_algo_method_tuning_gsu_gemm_only.groupedgemm_zero_n,extapi_groupedgemm(2),extapi_algo_method_grouped_gemm_tuning_wgm. Theirgpu_arch_exclude: '1[1-2]\d{2}'was meant for consumer gfx11/gfx12 but also matched gfx1250. It's now'11\d\d|12[^5]\d', the form other cases in the file already use._94x/_gfx942variants; gfx950 subtile FP4, tf32x1 and custom-kernel cases; rocRoller (gfx1250 builds skip it).size_t_stagger_incs_overflow_f32_ROCM31230, which the first gfx1250-strict run found has no solution); fp32 and f32-output bias gradients; MX cases using FP6 or fp16 output; f32 with fast bf16 compute in TT. These, and the gfx950 custom-kernel case, carry a# Not on gfx1250:comment. A case with no solution fails rather than skips, so enabling them would only add red.1250matches both device names. Strict-only failures are quarantined withknown_bug_platforms: "gfx1250-strict", which matches exactly.TESTING.mddocuments this, plus the consumer-exclusion idiom.gpu_archandgpu_arch_excludewiden from 16 to 32 bytes, since9(0a|42|50)|1250doesn't fit in 16.hipblaslt_gentest.pynow refuses any char-array value that doesn't fit with its terminator. Before, ctypes silently stored a value that exactly filled the field with no NUL.No library code changes.
Work tracking
Tests
Paths:
projects/hipblaslt/clients/tests/data/matmul_gtest.yaml,smoke_gtest.yaml: arch patterns and# Not on gfx1250:commentsprojects/hipblaslt/clients/tests/data/hipblaslt_common.yaml,clients/common/include/hipblaslt_arguments.hpp: field widthsprojects/hipblaslt/clients/tests/hipblaslt_gentest.py,clients/tests/test_hipblaslt_gentest.py: length check and its testsprojects/hipblaslt/TESTING.md: "Scoping test cases to architectures"Test results:
hipblaslt_gentest.py data/hipblaslt_gtest.yamlsucceeds. Output grows by exactly 32 bytes per record over 119,020 records, so the case count is unchanged. No existing char-array value reaches its limit.python -m unittest discover -s clients/tests -p 'test_hipblaslt_gentest.py': 4 pass. With the check weakened, 2 fail. These run by hand, likeclients/scripts/sdc_hunt; no CI job runs client Python tests today. The check itself runs in every build's test data generation.|1250or narrows an exclusion away from gfx1250, so the case set there is unchanged. Precheckin on this PR: pending.gfx1250 HW hipBLASLt + TensileLite Testsrun on this PR passed 17,373 hipBLASLt tests and failed 1:size_t_stagger_incs_overflow_f32_ROCM31230("NO solution found!"). That enable is reverted in c08a775 with the reason recorded. A rerun on that commit is pending. The run's one TensileLite failure was a 600 s timeout indecouple_pgr_tdm_fuse_gfx1250.yaml, which this PR doesn't touch and nightly develop runs have passed.Flags / guardrails
Adjacent tests considered
matmul_fast_check_mx_*already runs on gfx1250 and is the MX coverage there, so thematmul_mx_*cases that need FP6 or fp16-output MX stay off. test(hipblaslt): skip gfx1250-strict where the tests can't run #13464 narrows the bf8 rows inside those fast_check MX cases for the same reason; it touches different entries from this PR._94xi8 variants add activation epilogues. They stay off until i8 activation support on gfx1250 is confirmed; the plain i8 cases are enabled.Risk acceptance