Skip to content

test(miopen)[ALMIOPEN-2966]: exclude TheRock bump failures on gfx120X and Windows gfx1151 - #14882

Merged
aalabsi-amd merged 4 commits into
developfrom
users/aalabsi-amd/miopen-exclude-therock-bump-failures
Nov 3, 2026
Merged

aalabsi-amd merged 4 commits into
developfrom
users/aalabsi-amd/miopen-exclude-therock-bump-failures

Conversation

@aalabsi-amd

@aalabsi-amd aalabsi-amd commented Nov 2, 2026 •

Copy link
Copy Markdown
Contributor

JIRA-ID : ROCM-33858 ROCM-31323 ALMIOPEN-2966, ALMIOPEN-3129, ALMIOPEN-3130

Motivation

The rocm-libraries bump into TheRock went red with #13107 (ROCm/TheRock#9651, f4e24d5 → 8c41584; still red on #9680). #13107 added
Jenkins-only gtest families to the TheRock tiers, and TheRock tests gfx120X and Windows gfx1151, which the rocm-libraries PR gate doesn't
cover. Tracking issue: #14874.

Bump PR #9651 run 41113572856 failed 4 of 4 attempts. #9680 run
41405101044 also failed. Develop nightlies
41179300934 and
41283494663 were read across all attempts too.

# Failure Lane Evidence Action
1 GPU_ConvAIModelLoad_FP32.BundledFdeepModelsLoad: conv_ai_nd_heuristics_test.cpp:962 "no fdeep .tn.model files found" Linux + Windows gfx120X 10 of 10 executions on 8 hosts (#9651 attempts 1–4, #9680) exclude on gfx120X (ALMIOPEN-2966, already excluded on gfx110X/gfx115X)
2 GPU_SoftmaxFind20_FP16.softmaxBackward_log_{channel,instance}_mode_fp16 hang after "Testing solution functions..." → miopen_gtest_standard_gfx120X_suite ***Timeout 3600 s Linux gfx1201 (CS-RORDMZ-DT87, -DT81) channel: 3 of 3 executions on Linux gfx1201 (#9651, #9680). instance: hung on DT87 and took 27 s on DT81 in this PR's run 41516836287 once channel was excluded. Both pass in < 1.3 s on Linux gfx1200 and Windows gfx1201; BFP16/FP32 pass on Linux gfx1201 exclude GPU_SoftmaxFind20_FP16.* on gfx120X_linux (ALMIOPEN-3130). New with #13107 (GPU_SoftmaxFind20_*)
3 Smoke/CPU_DBSync_NONE.StaticFDBSync/11: "Db file does not exist … gfx120138.HIP.fdb.txt" Windows gfx120X, CS-RORDMZ-DT84 1 of 5 Windows gfx1201 executions. Only gfx120140 dbs ship; the other hosts pass exclude on gfx120X (#13416, same cause as StaticFDBSync/9 on gfx115X). Existed before #13107
4 CPU_PerfDb_NONE.{RamDb,PlainTextDb}_AllTests hang in the multiprocess read or write phase ("Waiting for test processes...") until the ctest or step timeout Windows gfx1151 6 hangs on 5 hosts since 10-31: nightly 41179300934 (RamDb), nightly 41283494663 (RamDb; PlainTextDb on attempts 1–3), #9680 (PlainTextDb). None in nightlies 40833587284 or 41018314721 exclude on gfx115X_windows (ALMIOPEN-3129)
5 Windows gfx1151 standard shard 3 ***Timeout 3600 s; Smoke/GPU_HipGraphExistTest_FP32.HipGraphExist/2 took 596 s (7.5 s elsewhere) #9680, windows-strix-halo-gpu-rocm-10 runner-specific slowness on top of #13107's larger standard tier not excluded: time budget, see below
6 Nightly Windows gfx1151 comprehensive shards hit the 120-min step cap (killed in Full/GPU_Pooling3d_FP16, Full/GPU_BnPeract_FP32, Full/GPU_Reduce_FP16) nightlies 10-31, 11-01 passed on rerun not excluded: time budget, see below
7 Linux gfx110X shard 2 ctest timeout (azure scale-set runner) #9651 attempt 1 passed on attempts 2–4 not excluded: flaky runner (TheRock#3845)
8 gfx90a CPU_PerfDb_NONE.MultiFileDb_AllTests "terminate called recursively" #9651 attempt 1 passed on attempts 2–4, seen once not excluded: one-off
9 gfx950 shards 1 and 3: no tests ran, Unexpected HTTP response: 429 #9651 attempt 1 infra not excluded

Technical Details

Only projects/miopen/test/gtest/test_categories.yaml changes. Each entry has the narrowest arch/OS scope that covers the failures, plus an
inline ticket or issue reference:

  • exclude_gpu_gfx120X: changed from the bare *gfx11_common_patterns alias to a list that also holds GPU_ConvAIModelLoad_FP32*
    (ALMIOPEN-2966) and Smoke/CPU_DBSync_NONE.StaticFDBSync/11 ([MIOpen] StaticFDBSync/11 fails on 56-CU gfx1201: no system database is shipped for that SKU #13416).
  • New exclude_gpu_gfx120X_linux: GPU_SoftmaxFind20_FP16.*, the two FP16 cases (ALMIOPEN-3130). The four BFP16/FP32
    SoftmaxFind20 cases keep running. It only hangs on gfx1201, but the exclusion covers all Linux gfx120X. TheRock takes the
    arch from AMDGPU_FAMILIES (gfx120X-all → gfx120X), so gfx1200 and gfx1201 runners share one suite. An exclude_gpu_gfx1201_*
    key would build a suite that no runner selects.
  • exclude_gpu_gfx115X_windows: CPU_PerfDb_NONE.RamDb_AllTests and CPU_PerfDb_NONE.PlainTextDb_AllTests (ALMIOPEN-3129). These
    two spawn child processes (DbMultiProcessReadTest, DbMultiProcessTest). MultiFileDb_AllTests doesn't, so it stays. Linux gfx115X
    already excludes CPU_PerfDb_NONE.* (ALMIOPEN-2964).

Jenkins gets no add_gtest_negative_filter change. It has no gfx120X or Windows stage, and these tests passed there before #13107
moved them into the TheRock tiers.

Not addressed here: slow Windows runners. After #13107, a standard shard runs about 6,550 cases (up from about 3,700), which takes
12–21 min on healthy Windows runners. A few runners are 5–7× slower per test and also take 8–11 min to start miopen_gtest.exe
(normally 0.1–0.2 min): windows-strix-halo-gpu-rocm-1/-2/-10 (gfx1151), navi4x-hw-32 (gfx1201), and, less so,
wshshark-strixhalo-16. Those runners hit or nearly hit the 60-min ctest limit on their own. Tests that start processes are worst hit:
MIOpenDriverRegressionHalf/0 (16.5 s normally, killed after 19.7 min) and HipGraphExist/2 (7.5 s, 596 s). That's a runner
problem, tracked in ROCM-33858. Excluding single tests or raising the timeout won't fix it.

Test Plan

  • python3 shared/ctest/parse_test_categories.py test_categories.yaml miopen_gtest …: rc 0, 25 suites. The generated
    miopen_gtest_standard_gfx120X_suite filter has the three gfx120X entries on Linux and two on Windows. The Windows
    miopen_gtest_standard_gfx115X_suite filter has the two PerfDb entries.
  • yaml_select.py diff --old git:origin/develop, using the 40,957 case names from the logs above, per arch/OS (every tier):
GPU / OS Removed (each tier)
gfx1201 Linux 4: ConvAIModelLoad, StaticFDBSync/11, SoftmaxFind20_FP16 ×2
gfx1201 Windows 2: ConvAIModelLoad, StaticFDBSync/11
gfx1151 Windows 2: CPU_PerfDb_NONE.{RamDb,PlainTextDb}_AllTests
gfx1151 Linux, gfx942 Linux, gfx1100 Windows 0
  • CI on this PR, with labels ci:gfx120X-all and ci:gfx1151. After it validates the gfx120X/gfx1151 lanes, drop
    ci:gfx120X-all so gfx94X can go green (see below).
  • After merge: the next TheRock bump should be green on gfx120X shards 1 and 3 and Windows gfx1151 shard 2.

Test Result

The selection diff above shows exactly the intended cases removed, with no change on other archs.

Run 41516836287 (commit 056ebc7, attempts 1 and 2). Every
exclusion took effect: in all 14 gfx120X/gfx1151 shard logs, the excluded cases are in the ctest gtest filter and never start.

Lane Result Notes
Linux gfx120X ×4 ✅ after rerun Attempt 1 shard 4 hung in GPU_SoftmaxFind20_FP16.softmaxBackward_log_instance_mode_fp16 on CS-RORDMZ-DT87 (gfx1201), the second FP16 case (ALMIOPEN-3130). The rerun on DT81 passed, but that case took 27.3 s. Fixed by 705e255, which excludes GPU_SoftmaxFind20_FP16.*.
Windows gfx120X ×4 ✅ Shard 1 on navi4x-hw-32 took 57.7 of 60 min (9.7-min startup, tests 5–7× slower). The other shards took 17–19 min.
Linux gfx1151 ×4 ✅ 10–14 min
Windows gfx1151 ×4 ✅ after rerun Attempt 1 shard 4 timed out on windows-strix-halo-gpu-rocm-1 (11.3-min startup, about 7× slower, killed in MIOpenDriverRegressionHalf/0). The rerun on wshsh3 finished in 12.5 min. No PerfDb hang.
Windows gfx110X ×4 ✅ 19–23 min, no selection change
Linux gfx94X ×4 ❌ 8 of 8 jobs Only the 17 Full/GPU_HipBLASLtGEMMTest_FP8 + 17 _BFP8 cases (hipblaslt_gemm.cpp:235, error 0.13–0.17 vs 0.07). Cause: #14375, not this PR (no change to the gfx94X selection). The ci:gfx120X-all label adds a gfx12 build. That job uploaded the shared miopen_{lib,run,test}_generic archives at 21:47:20, after gfx94X's (21:45:18), and the gfx94X tests downloaded them at 21:50. Per #14375, a host build whose targets include gfx12 sets MIOPEN_FP8_IEEE_EXPONENT_BIAS=1, which is wrong for gfx942. The rerun reused the same artifacts, so it failed the same way. Fix: #14627.

Submission Checklist

  • Each exclusion is scoped to arch/OS and cites a ticket or issue
  • YAML parses and the generator returns rc 0
  • Jira for the gfx1201 SoftmaxFind20 hang (ALMIOPEN-3130) and the Windows PerfDb multiprocess hang (ALMIOPEN-3129)
  • Remove ci:gfx120X-all (and ci:gfx1151) after validation, then get one green run with gfx94X
  • Infra ticket for the slow Windows runners (windows-strix-halo-gpu-rocm-1/-2/-10, navi4x-hw-32, wshshark-strixhalo-16): ROCM-33858
  • Re-enable each entry when its ticket closes
  • Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.

… and Windows gfx1151

Quarantine the MIOpen failures seen on TheRock bump PRs #9651/#9680
(issue #14874) after #13107 widened the TheRock tiers:

- gfx120X: GPU_ConvAIModelLoad_FP32* (no bundled fdeep models,
  ALMIOPEN-2966) and Smoke/CPU_DBSync_NONE.StaticFDBSync/11 (no
  gfx120138 system db on 56-CU gfx1201 hosts, #13416).
- Linux gfx120X: GPU_SoftmaxFind20_FP16.softmaxBackward_log_channel_mode_fp16
  hangs to the ctest timeout on gfx1201 hosts.
- Windows gfx115X: CPU_PerfDb_NONE.{RamDb,PlainTextDb}_AllTests hang in
  the multiprocess write phase.
@aalabsi-amd aalabsi-amd added project: miopen ci:gfx1151 Opt-in to gfx1151 builds/tests ci:gfx120X-all Opt-in to gfx120X-all builds/tests labels Nov 2, 2026
@therock-pr-bot

therock-pr-bot Bot commented Nov 2, 2026 •

Copy link
Copy Markdown

✅ All Checks Passed — Ready for Review

Check Status Details
📝 PR Description ✅ Pass —
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ✅ Pass PR does not contain code files — Unit Test auto-passed
🔎 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

therock-pr-bot Bot commented Nov 2, 2026

Copy link
Copy Markdown

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

@aalabsi-amd
aalabsi-amd marked this pull request as ready for review November 2, 2026 18:44
@aalabsi-amd
aalabsi-amd requested a review from a team as a code owner November 2, 2026 18:44

@JonathanLichtnerAMD JonathanLichtnerAMD left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks good. Exclusions are narrowly scoped, and the generator output for miopen_gtest_standard_gfx120X_suite contains the expected three entries on Linux. Three minor points:

  1. The Softmax hang is only observed on gfx1201 but is excluded for all Linux gfx120X. Scope it to gfx1201 if the key allows, or note that gfx1200 is excluded deliberately.
  2. StaticFDBSync/11 is an index-based name, so it will drift if the param list changes. Please make sure #13416 tracks removal.
  3. Please file Jira items for the Softmax and PerfDb hangs before merge and cite them in the YAML comments. They may be real bugs rather than CI-only issues.

Co-authored by Claude (AI).

… and Windows gfx1151

Quarantine the MIOpen failures seen on TheRock bump PRs #9651/#9680
(issue #14874) after #13107 widened the TheRock tiers:

- gfx120X: GPU_ConvAIModelLoad_FP32* (no bundled fdeep models,
  ALMIOPEN-2966) and Smoke/CPU_DBSync_NONE.StaticFDBSync/11 (no
  gfx120138 system db on 56-CU gfx1201 hosts, #13416).
- Linux gfx120X: GPU_SoftmaxFind20_FP16.softmaxBackward_log_channel_mode_fp16
  hangs to the ctest timeout on gfx1201 hosts.
- Windows gfx115X: CPU_PerfDb_NONE.{RamDb,PlainTextDb}_AllTests hang in
  the multiprocess write phase.
@aalabsi-amd

Copy link
Copy Markdown
Contributor Author

Looks good. Exclusions are narrowly scoped, and the generator output for miopen_gtest_standard_gfx120X_suite contains the expected three entries on Linux. Three minor points:

  1. The Softmax hang is only observed on gfx1201 but is excluded for all Linux gfx120X. Scope it to gfx1201 if the key allows, or note that gfx1200 is excluded deliberately.
  2. StaticFDBSync/11 is an index-based name, so it will drift if the param list changes. Please make sure [MIOpen] StaticFDBSync/11 fails on 56-CU gfx1201: no system database is shipped for that SKU #13416 tracks removal.
  3. Please file Jira items for the Softmax and PerfDb hangs before merge and cite them in the YAML comments. They may be real bugs rather than CI-only issues.

Co-authored by Claude (AI).

1- A YAML key can't limit the Softmax exclusion to gfx1201. TheRock's test_runner.py takes the arch from AMDGPU_FAMILIES |
| (gfx120X-all → gfx120X), so gfx1200 and gfx1201 runners both use the gfx120X suite. A YAML-only exclude_gpu_gfx1201_linux |
| key generates a separate gfx1201 suite that no runner selects, so the hang comes back on CS-RORDMZ-DT81/DT87. will keep on gfx120x

2- done and added to the issue
3- bugs created and will be mentioned in the yaml

@aalabsi-amd
aalabsi-amd enabled auto-merge (squash) November 2, 2026 21:34
… Linux gfx120X

GPU_SoftmaxFind20_FP16.softmaxBackward_log_instance_mode_fp16 hung on CS-RORDMZ-DT87
(PR #14882 run 41516836287 attempt 1) once channel_mode was excluded.
@assistant-librarian

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

yassinsolim added a commit that referenced this pull request Nov 4, 2026
…d system db (#14931)

ISSUE ID : #13416

## Motivation

`Smoke/CPU_DBSync_NONE.StaticFDBSync/11` (gfx1201) fails on gfx1201
devices whose CU count has no shipped system db. Only `gfx120140` (64
CU) ships. A 56-CU gfx1201 looks for `gfx120138.HIP.fdb.txt` and the
test stops at the file-exists assert in `SetupPaths` ("Db file does not
exist") before it checks any entry. Windows rockrel gfx120X-all
nightlies hit this on a 56-CU runner. This is the gfx1201 half of
#13416. The open PR #13179 uses the same mechanism for gfx1101.

## Technical Details

- `test/gtest/db_sync.cpp`: the gfx1201 `StaticFDBSync` param changes
from `0` to `32`. `Handle::GetMaxComputeUnits` reports WGPs on gfx12xx,
so 32 is the 64-CU part and maps to `gfx120140`. The existing check in
`StaticFDBSync` (`GetSysDbSelectionCu(arch, real_cu) != num_cu`) then
skips every other gfx1201 CU count. No new skip logic is added.
- The CU-mismatch skip now says why it skipped (device CU count and the
CU count whose db it selects). It used to print nothing.
- The param index stays 11, so the gtest name does not change.
- No CHANGELOG entry, this is test-only.

What could break or is left out:
- gfx1201 parts other than 64 CU (for example 56 and 48 CU) no longer
run `StaticFDBSync`. At runtime MIOpen still loads the closest shipped
database for them, and this test no longer checks that data. A 56-CU
check done earlier on this issue (emulated CU count) showed 83,001
Winograd entries from `gfx120140` that are invalid at 56 CU. Shipping a
`gfx120138` pair tuned on a 56-CU part is still the fix for that.
- gfx1100 (70, 80 and 84 CU), gfx1151 and gfx1200 params are unchanged,
so the same gap on those arches is not addressed here.
- #14882 excludes `Smoke/CPU_DBSync_NONE.StaticFDBSync/11` from TheRock
gfx120X CI. Once this and #14882 are in and TheRock picks up this
change, that exclusion can be dropped.

## Test Plan

Built `test_db_sync` for gfx1201 from this branch (CK for gfx1201 plus
MIOpen, ROCm 7.14.1 container, Ubuntu 24.04), and a second binary with
only `db_sync.cpp` reverted to develop as a control. No 56-CU gfx1201
was available, so the CU count was emulated with the rocjitsu KMD
interposer (same runner config as `run_dbsync_rocjitsu.py`,
`cu_correct_config` to 56, 48 and 64 CU).

1. 56-CU and 48-CU emulated gfx1201 with this branch: `StaticFDBSync/11`
must skip.
2. 56-CU emulated gfx1201 with the develop binary: expected to fail on
the missing file.
3. 64-CU emulated gfx1201 with this branch: `StaticFDBSync/11` must
still run the whole database check.
4. Real gfx1201 (Radeon AI PRO R9700, 64 CU) with this branch,
`--gtest_filter='*StaticFDBSync*'`: only `/11` runs, the other 11 params
skip.

## Test Result

1. Skipped at 56 CU and at 48 CU. With the filter `*StaticFDBSync*` at
56 CU all 12 params skip. The message printed is: `Skipping gfx1201
param 32: this device reports 28 compute units and selects the system db
for 28, which differs from this param.`
2. Develop binary at 56 CU: `Db file does not exist
".../gfx120138.HIP.fdb.txt"` and `[ FAILED ]
Smoke/CPU_DBSync_NONE.StaticFDBSync/11, where GetParam() = ("gfx1201",
0)`. This is the failure from the issue.
3. 64 CU emulated: `Handle CU count: 32 Parameter Value: 32`, reads
`gfx120140.HIP.fdb.txt`, `[ OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/11
(552657 ms)`, 1 test passed.
4. Real R9700: `Handle CU count: 32 Parameter Value: 32`, reads
`gfx120140.HIP.fdb.txt`, `[ OK ] StaticFDBSync/11 (11882 ms)`, 1 passed
and 11 skipped.

Not run: a real 56-CU gfx1201 (the skip at 56 CU is shown with an
emulated CU count), Windows, and the Jenkins or TheRock CI jobs.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
cycheng pushed a commit that referenced this pull request Nov 6, 2026
…d system db (#14931)

ISSUE ID : #13416

## Motivation

`Smoke/CPU_DBSync_NONE.StaticFDBSync/11` (gfx1201) fails on gfx1201
devices whose CU count has no shipped system db. Only `gfx120140` (64
CU) ships. A 56-CU gfx1201 looks for `gfx120138.HIP.fdb.txt` and the
test stops at the file-exists assert in `SetupPaths` ("Db file does not
exist") before it checks any entry. Windows rockrel gfx120X-all
nightlies hit this on a 56-CU runner. This is the gfx1201 half of
#13416. The open PR #13179 uses the same mechanism for gfx1101.

## Technical Details

- `test/gtest/db_sync.cpp`: the gfx1201 `StaticFDBSync` param changes
from `0` to `32`. `Handle::GetMaxComputeUnits` reports WGPs on gfx12xx,
so 32 is the 64-CU part and maps to `gfx120140`. The existing check in
`StaticFDBSync` (`GetSysDbSelectionCu(arch, real_cu) != num_cu`) then
skips every other gfx1201 CU count. No new skip logic is added.
- The CU-mismatch skip now says why it skipped (device CU count and the
CU count whose db it selects). It used to print nothing.
- The param index stays 11, so the gtest name does not change.
- No CHANGELOG entry, this is test-only.

What could break or is left out:
- gfx1201 parts other than 64 CU (for example 56 and 48 CU) no longer
run `StaticFDBSync`. At runtime MIOpen still loads the closest shipped
database for them, and this test no longer checks that data. A 56-CU
check done earlier on this issue (emulated CU count) showed 83,001
Winograd entries from `gfx120140` that are invalid at 56 CU. Shipping a
`gfx120138` pair tuned on a 56-CU part is still the fix for that.
- gfx1100 (70, 80 and 84 CU), gfx1151 and gfx1200 params are unchanged,
so the same gap on those arches is not addressed here.
- #14882 excludes `Smoke/CPU_DBSync_NONE.StaticFDBSync/11` from TheRock
gfx120X CI. Once this and #14882 are in and TheRock picks up this
change, that exclusion can be dropped.

## Test Plan

Built `test_db_sync` for gfx1201 from this branch (CK for gfx1201 plus
MIOpen, ROCm 7.14.1 container, Ubuntu 24.04), and a second binary with
only `db_sync.cpp` reverted to develop as a control. No 56-CU gfx1201
was available, so the CU count was emulated with the rocjitsu KMD
interposer (same runner config as `run_dbsync_rocjitsu.py`,
`cu_correct_config` to 56, 48 and 64 CU).

1. 56-CU and 48-CU emulated gfx1201 with this branch: `StaticFDBSync/11`
must skip.
2. 56-CU emulated gfx1201 with the develop binary: expected to fail on
the missing file.
3. 64-CU emulated gfx1201 with this branch: `StaticFDBSync/11` must
still run the whole database check.
4. Real gfx1201 (Radeon AI PRO R9700, 64 CU) with this branch,
`--gtest_filter='*StaticFDBSync*'`: only `/11` runs, the other 11 params
skip.

## Test Result

1. Skipped at 56 CU and at 48 CU. With the filter `*StaticFDBSync*` at
56 CU all 12 params skip. The message printed is: `Skipping gfx1201
param 32: this device reports 28 compute units and selects the system db
for 28, which differs from this param.`
2. Develop binary at 56 CU: `Db file does not exist
".../gfx120138.HIP.fdb.txt"` and `[ FAILED ]
Smoke/CPU_DBSync_NONE.StaticFDBSync/11, where GetParam() = ("gfx1201",
0)`. This is the failure from the issue.
3. 64 CU emulated: `Handle CU count: 32 Parameter Value: 32`, reads
`gfx120140.HIP.fdb.txt`, `[ OK ] Smoke/CPU_DBSync_NONE.StaticFDBSync/11
(552657 ms)`, 1 test passed.
4. Real R9700: `Handle CU count: 32 Parameter Value: 32`, reads
`gfx120140.HIP.fdb.txt`, `[ OK ] StaticFDBSync/11 (11882 ms)`, 1 passed
and 11 skipped.

Not run: a real 56-CU gfx1201 (the skip at 56 CU is shown with an
emulated CU count), Windows, and the Jenkins or TheRock CI jobs.

## Submission Checklist

- [x] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
aalabsi-amd added a commit that referenced this pull request Nov 6, 2026
…14169 (#15031)

JIRA ID : ALMIOPEN-1926

## Motivation

#14169 (merged 2026-10-26, `cec66d06`) fixed the data race in
`RNNDescriptor::RNNForwardMS` that made bias-enabled
`Full/GPU_LstmMSRnn_FP32.TestFloat` cases fail intermittently. Every
quarantine of this family still in the tree predates that fix:

| Where | Pattern | Added by | Data it was based on |
|---|---|---|---|
| Jenkins `MIOPEN_TEST_GFX94X` negative filter |
`Full/GPU_LstmMSRnn_FP32*` | #13107 | Jenkins gfx942, 2026-08-24..10-23
|
| `comprehensive.exclude`, `full.exclude` |
`Full/GPU_LstmMSRnn_FP32.TestFloat/1095` | #13107 | Jenkins per-index
filter |
| `exclude_gpu_gfx94X` | `Full/GPU_LstmMSRnn_FP32*` | #13107 | Jenkins
gfx942 (above) |
| `exclude_gpu_gfx110X` | `Full/GPU_LstmMSRnn_FP32*` | #13107 | Windows
gfx110X runs of 10-15/10-16 |
| `exclude_gpu_gfx115X` | `Full/GPU_LstmMSRnn_FP32*` | #13107 | Windows
gfx1151 run of 10-20 |

The #14169 author planned this cleanup as a follow-up once the fix was
confirmed. While these quarantines stay, a regression in the
LSTM multi-stream forward pass (reachable from `miopenRNNForward`,
`miopenRNNForwardTraining` and `miopenRNNForwardInference`) goes
untested on gfx942, gfx110X and gfx1151.

## Technical Details

- `projects/miopen/test/gtest/test_categories.yaml`:
  - remove `/1095` from the comprehensive and full excludes;
- remove `Full/GPU_LstmMSRnn_FP32*` from `exclude_gpu_gfx110X` and
`exclude_gpu_gfx115X`;
- remove the `exclude_gpu_gfx94X` key. Its only pattern was this family,
so gfx94X goes back to the base suites (`-LE ex_gpu`),
    as before #13107. That's 21 suites instead of 25.
- `projects/miopen/test/gtest/CMakeLists.txt`: remove the Jenkins
`MIOPEN_TEST_GFX94X` block, which only held this filter.
- **Kept:** `exclude_gpu_gfx115X_windows`
`Full/GPU_GRU_Base_FP16…biasMode_1_*`. GRU doesn't use the multi-stream
LSTM forward that
  #14169 changed, so the fix doesn't cover that flake.
- **Kept:** `standard.exclude` `Full/GPU_LstmMSRnn_*`. That's tier
placement (the Full/ cases run in comprehensive), not a quarantine.

## Test Plan

1. `parse_test_categories.py`: rc 0, 21 suites.
2. `yaml_select.py diff --old git:origin/develop`, using the 69,069 case
names from `miopen_gtest --gtest_list_tests` of the post-fix build:

| GPU / OS | standard | comprehensive / full |
|---|---|---|
| gfx942 Linux | 0 | +1,536 `Full/GPU_LstmMSRnn_FP32` |
| gfx1100 Windows | 0 | +1,536 |
| gfx1151 Windows / Linux | 0 | +1,536 |
| gfx1201 Linux, gfx950 Linux | 0 | +1 (`/1095`; they already ran the
rest) |

3. **Local A/B stress on gfx942 (MI300X).** Same method as #14169. The
pre-fix build is TheRock CI artifacts of rl run 39358709264
(2026-10-20); the post-fix build is rl run 41582846076 (develop
`9953cb5` + PR #14882). Two modes: (A) one process with
`GPU_MAX_HW_QUEUES=16`; (B) 4 gtest shards sharing one GPU, as Jenkins
does. Rounds alternate pre/post, each with a fresh user db
   and kernel cache. Script: `/scratch/users/aalabsi/tmp/lstm_ab/ab.sh`.
4. **CI on this PR**, labels `test:miopen` + `test_filter:comprehensive`
+ `ci:gfx1151`. This runs
comprehensive (including all 1,536 `Full/GPU_LstmMSRnn_FP32` cases) on
Linux gfx94X, Windows gfx110X and Windows/Linux gfx1151.
   These are the lanes that failed before. Read every attempt.
5. **Jenkins PR**: the gfx942 stages run the family again now that the
filter is gone.

## Test Result

**Local A/B stress, gfx942 (MI300X, banff-cyxtera-s70-4)**, Slurm jobs
68262769 (3 rounds) and 68262798 (20 rounds),
23 rounds per build per mode, pre/post alternating. Each round runs all
1,536 cases: 1,152 execute and 384 skip by design at
`rnn_seq_api.cpp:61` in both builds.

| Build | Mode A: 1 process, `GPU_MAX_HW_QUEUES=16` | Mode B: 4 shards
on 1 GPU (Jenkins-like) |
|---|---|---|
| pre-fix (rl run 39358709264, 2026-10-20) | 0 failures / 26,496 | **150
failures / 26,496 (0.57%), in 23 of 23 rounds, 113 distinct indices**,
all at `rnn_seq_api.hpp:272` (the ALMIOPEN-1926 signature) |
| post-fix (rl run 41582846076, develop `9953cb5`) | 0 / 26,496 | **0 /
26,496, in 0 of 23 rounds** |

At the pre-fix rate, 26,496 post-fix executions would give about 150
failures. Seeing 0 by chance is effectively impossible
(p ≈ e^-150). Mode A didn't reproduce on this machine for either build;
#14169 also needed mode B or many passes. A round takes
77–81 s for both builds, so the fix costs no measurable test time. Logs:
`/scratch/users/aalabsi/tmp/lstm_ab/results-68262769/`,
`…/results-68262798/`.

**Jenkins develop, gfx942** (QE dashboard): 20 failures / 11,719
executions from 2026-09-01 to 10-26; 0 / 1,144 from 10-27 to 10-30,
after which #13107 filtered the whole family. gfx90a: 0 / 11,615 before
and 0 / 1,872 after; it was never quarantined.

**CI on this PR:** pending. Windows gfx110X and gfx1151 are the open
question: no Windows run of this family exists after the fix,
and the stress test above covers gfx942 only.

## Submission Checklist

- [x] Every removed entry cites ALMIOPEN-1926 and is covered by #14169
- [x] YAML parses; selection diff shows only the intended cases
- [x] gfx942 A/B: pre-fix fails under stress (150 / 26,496), post-fix
doesn't (0 / 26,496)
- [ ] PR CI comprehensive: Windows gfx110X / gfx1151 and Linux gfx94X
run `Full/GPU_LstmMSRnn_FP32` with 0 failures in every attempt
- [ ] Jenkins gfx942 PR stages green with the filter removed
- [ ] Remove the `test_filter:comprehensive` label before merge
- [ ] Close ALMIOPEN-1926, or narrow it to the GRU FP16 flake
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

ci:gfx120X-all Opt-in to gfx120X-all builds/tests ci:gfx1151 Opt-in to gfx1151 builds/tests organization: ROCm project: miopen

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants