test(hip-kernel-provider): port mlops engine integration tests#9255
Draft
adickin-amd wants to merge 4 commits into
Draft
test(hip-kernel-provider): port mlops engine integration tests#9255adickin-amd wants to merge 4 commits into
adickin-amd wants to merge 4 commits into
Conversation
Ports hip-kernel-provider's mlops-engine-local Batchnorm/RMSnorm/Layernorm integration tests into the shared dnn-providers/integration-tests suite, wires hip-kernel-provider onto add_external_integration_test_target(), and adopts miopen-provider's YAML-driven quick/standard/comprehensive/full CTest tier framework, mirroring the pattern used for hipblaslt-provider (PR #9205). Common suite (dnn-providers/integration-tests): - Batchnorm and RMSNorm-forward already existed generically and are a strict superset of the mlops-local coverage (1D support, full/regression variants, golden references, parametrized activation backward modes) - no new code needed there. - Add src/integration_tests/rmsnorm/IntegrationGpuRMSNormBackward.cpp, reusing the existing RMSNormTestCase/getRMSNormTestCases() generator (bias/isTraining are forward-only concerns rmsnorm_backward() ignores). - Add src/common/LayernormCommon.hpp and src/integration_tests/layernorm/IntegrationGpuLayernormForward.cpp (net new - no layernorm coverage existed anywhere in the common suite yet). - Extend harness/IntegrationGraphVerificationHarness.hpp's tolerance dispatch with RMSNormBackwardNode and LayerNormNode entries (both reuse existing hipdnn_test_sdk tolerance functions - no TestTolerances.hpp changes needed). hip-kernel-provider: - Remove the now-ported src/integration_tests/hip_mlops_engine/ tree (Batchnorm/, RMSnorm/, Layernorm/, Common/ActivationCommon.hpp) and its CMakeLists.txt wiring. The provider-local IntegrationGraphVerificationHarness.hpp and asm_sdpa_engine/ tests are untouched - the harness is still shared with the (out of scope) ASM SDPA engine. - Add config/HIP_MLOPS_ENGINE.toml (template only - no known skips or tolerance overrides yet). - Add add_external_integration_test_target() wiring in src/CMakeLists.txt (ENGINE_NAME HIP_MLOPS_ENGINE, PLUGIN_TARGET hip_kernel_provider) plus the ROCM_LIBRARIES_ROOT guard in the top-level CMakeLists.txt. - Add test_categories.yaml, test_categories_external.yaml, and test_categories_integration.yaml, adopting miopen-provider's tiering verbatim (hip-kernel-provider had no YAML-based CTest categorization before this). Verified: superbuild of the hip-kernel-provider preset builds clean under -Werror. hipdnn_integration_tests --test-engine HIP_MLOPS_ENGINE --test-config config/HIP_MLOPS_ENGINE.toml --gtest_filter="*Batchnorm*:*RMSNorm*:*Layernorm*" runs 2388 cases, 1690 pass, 698 skip (golden-ref CPU-only cases plus the intentional inference-mode mixed-precision layernorm skip carried over from the provider-local suite), 0 failures.
hip-kernel-provider has more than one engine (ASM_SDPA_ENGINE, HIP_MLOPS_ENGINE); rename test_categories*.yaml to HIP_MLOPS_ENGINE_test_categories*.yaml so a future per-engine external-integration-check target (e.g. for ASM_SDPA_ENGINE) doesn't collide with or overload these filenames. Matches the existing config/HIP_MLOPS_ENGINE.toml naming convention. No behavior change: reran the HIP_MLOPS_ENGINE external integration suite post-rename (2388 tests, 1690 pass, 698 skip, 0 fail - identical to pre-rename).
Contributor
Author
|
Note: not ready to review, just putting up to see if it builds on CI and what needs to change on the yaml filters. |
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
…cks conv Running the full (unfiltered) hipdnn_integration_tests suite against HIP_MLOPS_ENGINE surfaced 18 real failures, all in IntegrationGpuConvFwdSerializeRoundTripFp32: unlike every other test, it calls graphObj.build() directly instead of going through the harness's verifyGraph(), so it never checks engine support before asserting build() succeeds - a no-conv engine like HIP_MLOPS_ENGINE fails the ASSERT_EQ instead of skipping. Extracts the existing engine-support-check block out of IntegrationGraphVerificationHarness::verifyGraph() into a new checkEngineSupportOrSkip() helper (identical logic, no behavior change for existing callers) and calls it from the conv serialize round-trip test before its manual graphObj.build(). Verified: full unfiltered hipdnn_integration_tests run against HIP_MLOPS_ENGINE now reports 0 failures (3939 tests, 1692 passed, 2247 skipped - the 18 conv cases moved from FAILED to SKIPPED; passed count is unchanged, confirming no regression for MIOPEN_ENGINE/HIPBLASLT_ENGINE-style callers).
…_kernel_provider_mlops_integration_test_port # Conflicts: # dnn-providers/integration-tests/src/harness/IntegrationGraphVerificationHarness.hpp
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.
Summary
Ports hip-kernel-provider's mlops-engine-local Batchnorm/RMSnorm/Layernorm integration tests into the shared
dnn-providers/integration-testssuite and wires the provider ontoadd_external_integration_test_target()plus the YAML-driven quick/standard/comprehensive/full CTest tier framework, mirroring the pattern already merged for miopen-provider. JIRA ID : ALMIOPEN-2283This covers the mlops-engine half of ALMIOPEN-2283 only; the story also asks for the ASM SDPA engine and rocKE is explicitly deferred there, but this PR intentionally leaves both untouched (rocKE per the story, ASM SDPA per direction from the assignee) to keep the change scoped to one engine.
Running the full (unfiltered) common suite against HIP_MLOPS_ENGINE for the first time - exactly the "does the external test cmake stuff all work as expected" unknown the story calls out - surfaced one real, pre-existing bug unrelated to mlops itself:
IntegrationGpuConvFwdSerializeRoundTripFp32bypasses the harness's usual engine-support check and assertsgraphObj.build()succeeds unconditionally, so it fails outright (instead of skipping) against any engine that doesn't implement conv. Fixed by extracting the existing check into a reusablecheckEngineSupportOrSkip()harness helper and calling it from that test too.Risk Assessment
Medium risk. The diff itself is test/build-only - no production or kernel code changed - and is thoroughly covered locally (full unfiltered suite: 3939 tests, 0 failures on gfx1151). However, it activates previously-inactive generic integration coverage (Batchnorm, RMSNorm) against HIP_MLOPS_ENGINE through a new entrypoint and adds net-new generic coverage (RMSNorm backward, Layernorm forward), so per policy it needs a full multi-arch sweep before merge - this exact engine/op family has a recent history of arch-specific failures (#9139, #9151, both already fixed upstream and confirmed passing on this branch) that a sweep could plausibly resurface on unswept targets, and running the full suite here already surfaced one real (now-fixed) common-suite gap.
ASIC Coverage
Full multi-arch sweep required. This activates generic integration test cases (Batchnorm, RMSNorm, Layernorm) in a new provider lane (HIP_MLOPS_ENGINE via
hip-kernel-provider-external-integration-check) and adds two net-new generic test files; a newly activated or added generic case can fail on any target, so standard PR CI (routine families: gfx94X, gfx950 on Linux, gfx1151 on Windows) is not sufficient on its own. The full supported family set (gfx94X-all, gfx950-all, gfx110X-all, gfx120X-all, plus gfx1151) should pass before merge.Testing Summary
hip-kernel-providerpreset builds clean under-Werror(699/699 targets).hipdnn_integration_testssuite unfiltered, pinned toHIP_MLOPS_ENGINEvia--test-engine, on local gfx1151 hardware: 3939 tests, 1692 passed, 2247 skipped (ops HIP_MLOPS_ENGINE doesn't implement + golden-ref CPU-only cases + the intentional inference-mode mixed-precision layernorm skip), 0 failed.Testing Checklist
cmake --build build --target hipdnn_integration_tests hip_kernel_provider_integration_tests(hip-kernel-provider preset,-Werror) - Status: Passedhipdnn_integration_tests --test-engine HIP_MLOPS_ENGINE --test-config config/HIP_MLOPS_ENGINE.toml- ASICs: gfx1151 - Status: Passed (3939 ran, 1692 passed, 2247 skipped, 0 failed)Technical Changes
dnn-providers/integration-tests: Batchnorm and RMSNorm-forward already existed generically and are a strict superset of the mlops-local coverage (1D support, full/regression tiers, golden references, parametrized backward-activation modes); no new code needed for those.dnn-providers/integration-tests: addsintegration_tests/rmsnorm/IntegrationGpuRMSNormBackward.cpp, reusing the existingRMSNormTestCase/getRMSNormTestCases()generator (bias/isTraining are forward-only concerns thatrmsnorm_backward()ignores).dnn-providers/integration-tests: addscommon/LayernormCommon.hppandintegration_tests/layernorm/IntegrationGpuLayernormForward.cpp- net new, no layernorm coverage existed anywhere in the common suite before this.dnn-providers/integration-tests: extends the harness's tolerance dispatch (IntegrationGraphVerificationHarness.hpp) withRMSNormBackwardNode/LayerNormNodeentries, both reusing existinghipdnn_test_sdktolerance functions.dnn-providers/integration-tests: extracts the harness's existing engine-support-check block out ofverifyGraph()into a reusablecheckEngineSupportOrSkip()helper (no behavior change for existing callers) and calls it fromIntegrationGpuConvForwardSerializeRoundTrip.cpp, which was bypassing that check entirely and failing (rather than skipping) against engines that don't implement conv, such as HIP_MLOPS_ENGINE.dnn-providers/hip-kernel-provider: removes the now-portedsrc/integration_tests/hip_mlops_engine/tree (Batchnorm/, RMSnorm/, Layernorm/, localActivationCommon.hpp) and its CMake wiring. The provider-localIntegrationGraphVerificationHarness.hppandasm_sdpa_engine/tests are untouched, since that harness is still shared with the (out of scope) ASM SDPA engine.dnn-providers/hip-kernel-provider: addsconfig/HIP_MLOPS_ENGINE.toml(template only, no known skips or tolerance overrides yet) and wiresadd_external_integration_test_target()(ENGINE_NAME HIP_MLOPS_ENGINE,PLUGIN_TARGET hip_kernel_provider).dnn-providers/hip-kernel-provider: addsHIP_MLOPS_ENGINE_test_categories.yaml,HIP_MLOPS_ENGINE_test_categories_external.yaml, andHIP_MLOPS_ENGINE_test_categories_integration.yaml, adopting miopen-provider's tiering. Files are engine-name-prefixed (rather than the generictest_categories*.yamlmiopen/hipblaslt use) since hip-kernel-provider has more than one engine (ASM_SDPA_ENGINE, HIP_MLOPS_ENGINE) and a future per-engine external-integration-check target shouldn't collide with these names.