Repository navigation
test(miopen)[ALMIOPEN-1926]: re-enable Full/GPU_LstmMSRnn_FP32 after #14169 - #15031
Conversation
…14169 #14169 (merged 2026-10-26) fixed the data race in RNNDescriptor::RNNForwardMS that made bias-enabled Full/GPU_LstmMSRnn_FP32.TestFloat cases fail intermittently. Remove the quarantines that cite ALMIOPEN-1926 for this family: - test_categories.yaml: Full/GPU_LstmMSRnn_FP32.TestFloat/1095 in the comprehensive and full excludes; Full/GPU_LstmMSRnn_FP32* in exclude_gpu_gfx110X and exclude_gpu_gfx115X; the exclude_gpu_gfx94X key, whose only pattern was this family. - CMakeLists.txt: the Jenkins MIOPEN_TEST_GFX94X negative filter. The Windows gfx115X GRU FP16 entry stays: GRU does not use the multi-stream LSTM forward that #14169 fixed.
✅ 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. |
JonathanLichtnerAMD
left a comment
There was a problem hiding this comment.
The yaml and CMake cleanup is correct and complete. MIOPEN_TEST_GFX94X is still used elsewhere, so removing only the filter block is right. The kept GRU FP16 and standard-tier entries are also right.
Needs a follow-up when merging with develop: #15021 added projects/miopen/AGENTS.md, whose example ctest ... -L '^ex_gpu_gfx94X$' uses the label this PR deletes. After merge that command fails with --no-tests=error. Please update it in this PR once you rebase or merge develop, or in a follow-up. Suggested change (untested; ex_gpu_gfx950 still exists at test_categories.yaml:753):
-# test_categories.yaml, often a family name (gfx942 -> ex_gpu_gfx94X).
-ctest ... -L '^quick$' -L '^ex_gpu_gfx94X$'
+# test_categories.yaml, often a family name (gfx950 -> ex_gpu_gfx950).
+ctest ... -L '^quick$' -L '^ex_gpu_gfx950$'Request: read every Windows gfx110X and gfx1151 attempt of Full/GPU_LstmMSRnn_FP32, because the stress evidence covers gfx942 only.
No blockers. 0 suggestion blocks are included, since the AGENTS.md line is not in this PR's diff.
Co-authored by Claude (AI).
…miopen-reenable-lstmmsrnn
…l that still exists This PR removes exclude_gpu_gfx94X, so ctest -L '^ex_gpu_gfx94X$' would select nothing and fail under --no-tests=error. Use gfx1100 -> ex_gpu_gfx110X, which keeps the family-name example.
JIRA ID : ALMIOPEN-1926
Motivation
#14169 (merged 2026-10-26,
cec66d06) fixed the data race inRNNDescriptor::RNNForwardMSthat made bias-enabledFull/GPU_LstmMSRnn_FP32.TestFloatcases fail intermittently. Every quarantine of this family still in the tree predates that fix:MIOPEN_TEST_GFX94Xnegative filterFull/GPU_LstmMSRnn_FP32*comprehensive.exclude,full.excludeFull/GPU_LstmMSRnn_FP32.TestFloat/1095exclude_gpu_gfx94XFull/GPU_LstmMSRnn_FP32*exclude_gpu_gfx110XFull/GPU_LstmMSRnn_FP32*exclude_gpu_gfx115XFull/GPU_LstmMSRnn_FP32*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,miopenRNNForwardTrainingandmiopenRNNForwardInference) goesuntested on gfx942, gfx110X and gfx1151.
Technical Details
projects/miopen/test/gtest/test_categories.yaml:/1095from the comprehensive and full excludes;Full/GPU_LstmMSRnn_FP32*fromexclude_gpu_gfx110Xandexclude_gpu_gfx115X;exclude_gpu_gfx94Xkey. Its only pattern was this family, so gfx94X goes back to the base suites (-LE ex_gpu),as before test(miopen): add Jenkins-only gtest families to test_categories.yaml… #13107. That's 21 suites instead of 25.
projects/miopen/test/gtest/CMakeLists.txt: remove the JenkinsMIOPEN_TEST_GFX94Xblock, which only held this filter.exclude_gpu_gfx115X_windowsFull/GPU_GRU_Base_FP16…biasMode_1_*. GRU doesn't use the multi-stream LSTM forward thatfix(miopen): order LSTM multi-stream fwd bias add before GEMM #14169 changed, so the fix doesn't cover that flake.
standard.excludeFull/GPU_LstmMSRnn_*. That's tier placement (the Full/ cases run in comprehensive), not a quarantine.Test Plan
parse_test_categories.py: rc 0, 21 suites.yaml_select.py diff --old git:origin/develop, using the 69,069 case names frommiopen_gtest --gtest_list_testsof the post-fix build:Full/GPU_LstmMSRnn_FP32/1095; they already ran the rest)(2026-10-20); the post-fix build is rl run 41582846076 (develop
9953cb5+ PR test(miopen)[ALMIOPEN-2966]: exclude TheRock bump failures on gfx120X and Windows gfx1151 #14882). Two modes: (A) one process withGPU_MAX_HW_QUEUES=16; (B) 4 gtest shards sharing one GPU, as Jenkins does. Rounds alternate pre/post, each with a fresh user dband kernel cache. Script:
/scratch/users/aalabsi/tmp/lstm_ab/ab.sh.test:miopen+test_filter:comprehensive+ci:gfx1151. This runscomprehensive (including all 1,536
Full/GPU_LstmMSRnn_FP32cases) on Linux gfx94X, Windows gfx110X and Windows/Linux gfx1151.These are the lanes that failed before. Read every attempt.
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:61in both builds.GPU_MAX_HW_QUEUES=16rnn_seq_api.hpp:272(the ALMIOPEN-1926 signature)9953cb5)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
Full/GPU_LstmMSRnn_FP32with 0 failures in every attempttest_filter:comprehensivelabel before merge