test(hipdnn): move the heaviest layernorm cases out of the nightly comprehensive tier - #12221
Draft
adickin-amd wants to merge 1 commit into
Draft
adickin-amd wants to merge 1 commit into
adickin-amd wants to merge 1 commit into
Conversation
hip-kernel-provider-external-integration_comprehensive_suite still exceeds the nightly job's 30-minute wall on every arch, so the nightly has no usable signal. Timings parsed from the 2026-09-16 gfx110X-all nightly (run 35048741357, job 104671445871) show it reached only 2957 of 6513 selected tests before the kill, with the cost concentrated in layernorm: full_Layernorm_Default 838.0 s (320 of 336 cases, killed inside) all other full_* bundles 300.0 s (1800 cases) Comprehensive/ layernorm 382.2 s (280 cases) Move both layernorm groups to the weekly full tier. No case is deleted or skipped. The comprehensive tier keeps full_* and excludes full_Layernorm_* by suite name, so the cheap full-tier batchnorm and convolution bundles keep their nightly coverage. Both Layernorm sweeps are excluded, not just Default: Variant2 is 144 further cases at ~0.4x Default's element load and was never reached in the failing run. The C++ 5D Comprehensive catalog is folded into Full, making the 5D tier boundary Quick / Standard (<=400K elements) / Full (>400K). Shape count is conserved: 36 Full shapes x 2 layouts x 7 backward fixtures = 504, exactly the previous 224 Full + 280 Comprehensive, and the 280 matches the count measured in the failing nightly. Because all three provider test_categories_integration.yaml files list Comprehensive/* and none list Full/*, this also moves those shapes from nightly to weekly on the miopen and hipblaslt lanes. That is intended: adding Full/* to their comprehensive tiers instead would pull the >5M element weekly shapes into their nightlies. Comprehensive/* is retained in the HIP_MLOPS pattern list as a no-op extension point so all three provider files keep the same shape. Follow-up to #11884; #11529 is still open and reproducing.
✅ 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. |
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
1 participant
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.
Summary
hip-kernel-provider-external-integration_comprehensive_suitestill exceeds thenightly job's 30-minute wall on every arch, so the nightly has no usable signal.
#11884 retiered part of this surface; the measured remainder is the full-tier
Layernorm bundles (which the comprehensive tier pulls in via
full_*) and theC++
Comprehensive/layernorm-backward shapes. This moves both to the weeklyfulltier. No case is deleted or skipped.JIRA ID : ROCM-31126
Follow-up to #11884; #11529 is still open and reproducing.
Timings parsed from the 2026-09-16
gfx110X-allnightly (run 35048741357, job104671445871),
which reached only 2957 of 6513 selected tests before the kill:
full_Layernorm_Defaultfull_*bundlesComprehensive/layernormStandard/+Quick/+Smoke/quick_*/standard_*bundlesRisk Assessment
Risk 3. No product code changes — the diff is a test shape catalog, gtest
instantiations, and one CTest category file. The blast radius is wider than the
one failing job, though: removing the
Comprehensivelayernorm instantiationsmoves those cases out of the nightly tier for the miopen and hipblaslt lanes
too, because all three provider
test_categories_integration.yamlfiles listComprehensive/*and none listFull/*. That is a deliberate cadence change(nightly to weekly) for those shapes across all providers, not a
hip-kernel-provider-only edit, and it is the main thing to review. The
alternative — adding
Full/*to the other two providers' comprehensive tiers topreserve nightly coverage — would pull the >5M-element weekly shapes into their
nightlies and break them instead. Case count is conserved end to end, so
residual risk is tier routing, not correctness.
ASIC Coverage
Passing PR CI is sufficient; no multi-arch sweep required. Nothing here can
change kernel selection, support surface, or default behavior on any target, and
no tier gains cases — every change moves a case to a heavier tier. The decisive
evidence is the next nightly on the affected families (Linux
gfx94X-dcgpu,gfx110X-all,gfx120X-all,gfx1151; Windowsgfx110X-all,gfx120X-all,gfx1151), since the failure this fixes only manifests in the nightlycomprehensive lane.
Testing Summary
hip-kernel-providersuperbuild, gfx1151/Windows, 869/869 targets.build tree and checked against the intended filter. Resolved comprehensive
filter is
quick_*:standard_*:full_*:Smoke/*:Quick/*:Standard/*:Comprehensive/*:Integration*-*DISABLED*:full_Layernorm_*.counts measured in the failing nightly (36 shapes x 2 layouts x 7 backward
fixtures = 504 = the previous 224 Full + 280 Comprehensive; the 280 matches the
nightly exactly).
Not verified locally: the integration-tests binary does not launch on this
workstation (exit 53, no output, even for
--gtest_help; under ctest itsurfaces as
0xc0000135). An untouched prior build fails identically, so this isa pre-existing local environment problem, not a property of this change. The
post-change wall-clock is therefore projected from the nightly log rather than
measured, and PR CI plus the next nightly are the real evidence.
Testing Checklist
cmake --build build --parallel 16(presethip-kernel-provider, gfx1151) - Status: Passedctest -R '^hip-kernel-provider-external-integration_comprehensive_suite$' -V -N- Status: Passedpython3 shared/ctest/parse_test_categories.py <yaml> hkp-ext /tmp- Status: Passedpre-commit run --files <4 changed files>- Status: PassedTechnical Changes
HIP_MLOPS_ENGINE_test_categories_integration.yaml: comprehensive tier keepsfull_*but excludesfull_Layernorm_*. Written as an explicitexclude:list because YAML cannot append to the
*common_excludesalias. Both Layernormsweeps are excluded, not just
Default:Variant2is 144 further cases at~0.4x
Default's element load and was never reached in the failing run.LayernormCommon.hpp:getLayernorm5DComprehensiveTestCases()removed; its 20shapes folded into
getLayernorm5DFullTestCases(). The 5D tier boundary is nowQuick / Standard (<=400K elements) / Full (>400K).
IntegrationGpuLayernormBackward.cpp,IntegrationGpuLayernormForward.cpp:the 7
INSTANTIATE_TEST_SUITE_P(Comprehensive, ...)blocks removed from each.Comprehensive/*is retained in the HIP_MLOPS pattern list as a no-opextension point so all three provider files keep the same shape.
Note for reviewers — M3
Nothing is disabled, skipped, or weakened: all 504 backward cases still run, and
full_Layernorm_*still runs in the weeklyfulltier. The honest framing isthat nightly coverage of those shapes drops to weekly, and the trigger was a red
nightly. The position here is that this corrects a tier mis-assignment —
weekly-cost work was scheduled nightly — rather than weakening tests. Flagging it
explicitly rather than leaving it implicit.
Known gate failure
The Libraries PR Bot
unit_testscheck is expected to fail. It requires any PRtouching
.cpp/.hppto also touch a file matchingtest_*,testing_*,*_test.*,*_tests.*,Test*,*_gtest.*, or**/test/gtest/**. All threecode files here are tests, but they are named
IntegrationGpu*.cpp/LayernormCommon.hppunderintegration-tests/src/, which matches none of thosepatterns. This is pre-existing: #11884 touched the same files with the same
non-matching names. The durable fix is a path pattern for hipDNN's integration
tests in
tools/libraries_pr_bot/policy.yml, which belongs in its own PR.