From f53d46fcbd6676ebf86a0cd666651235b170833e Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 09:00:25 -0600 Subject: [PATCH 01/16] perf(hipdnn): compare plan registry keys by index before one visit PlanRegistrySignatureKeyEqual and GpuPlanRegistrySignatureKeyEqual used std::visit(*this, a, b), which instantiates the visitor for every pair of alternatives (22x22 = 484 and 15x15 = 225). Compare the active alternative indices first, then visit only `a` and fetch the same alternative from `b` (22 and 15 instantiations). Semantics are unchanged: keys of different alternatives are unequal, keys of the same alternative compare with ==. Add TestGpuPlanRegistrySignatureKey, pinning same-type equal and hash-equal, same-type different-value unequal, cross-type unequal, and unordered_map lookup with the key hash/equal functors. --- .../detail/GpuPlanRegistrySignatureKey.hpp | 25 ++- .../integration-tests/tests/CMakeLists.txt | 1 + .../tests/TestGpuPlanRegistrySignatureKey.cpp | 165 ++++++++++++++++++ .../detail/PlanRegistrySignatureKey.hpp | 25 ++- 4 files changed, 188 insertions(+), 28 deletions(-) create mode 100644 dnn-providers/integration-tests/tests/TestGpuPlanRegistrySignatureKey.cpp diff --git a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp index 6e89b6f8a56d..fbd6db563437 100644 --- a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp +++ b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp @@ -4,6 +4,7 @@ #pragma once #include +#include #include #include "GpuBatchnormFwdInfSignatureKey.hpp" @@ -51,23 +52,19 @@ struct GpuPlanRegistrySignatureKeyHash struct GpuPlanRegistrySignatureKeyEqual { - template - bool operator()([[maybe_unused]] const T& a, [[maybe_unused]] const U& b) const - { - return false; - } - - template - bool operator()(const T& a, const T& b) const - { - return a == b; - } - - // NOLINTNEXTLINE(readability-redundant-casting) + // Compare the active alternatives' indices first, then visit only `a` and fetch the matching + // alternative from `b`. Visiting both variants would instantiate one visitor call per pair of + // alternatives (N*N) instead of N. bool operator()(const GpuPlanRegistrySignatureKey& a, const GpuPlanRegistrySignatureKey& b) const { - return std::visit(*this, a, b); + if(a.index() != b.index()) + { + return false; + } + + return std::visit( + [&b](const auto& x) { return x == std::get>(b); }, a); } }; diff --git a/dnn-providers/integration-tests/tests/CMakeLists.txt b/dnn-providers/integration-tests/tests/CMakeLists.txt index cdc7a9605e2d..dbadb3b5987b 100644 --- a/dnn-providers/integration-tests/tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/tests/CMakeLists.txt @@ -61,6 +61,7 @@ add_executable(hipdnn_integration_tests_unit_tests TestGpuBatchnormFwdInfPlan.cpp TestGpuBatchnormFwdInfVarianceSignatureKey.cpp TestGpuBatchnormFwdInfVariancePlan.cpp + TestGpuPlanRegistrySignatureKey.cpp TestReferenceGraphExecutorFactory.cpp TestBundleDiscovery.cpp TestGTestFilter.cpp diff --git a/dnn-providers/integration-tests/tests/TestGpuPlanRegistrySignatureKey.cpp b/dnn-providers/integration-tests/tests/TestGpuPlanRegistrySignatureKey.cpp new file mode 100644 index 000000000000..a5b362ed2ce1 --- /dev/null +++ b/dnn-providers/integration-tests/tests/TestGpuPlanRegistrySignatureKey.cpp @@ -0,0 +1,165 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include +#include +#include +#include + +#include +#include +#include + +#include "MatmulGraphTestUtils.hpp" +#include "harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp" + +using namespace hipdnn_data_sdk::utilities; +using namespace hipdnn_flatbuffers_sdk::data_objects; +using namespace hipdnn_integration_tests::test_utils; +using namespace hipdnn_integration_tests::gpu_graph_executor::detail; + +namespace +{ + +// Builds a matmul key from a real flatbuffer node and tensor map. +GpuPlanRegistrySignatureKey + makeMatmulKey(DataType aType, DataType bType, DataType cType, DataType computeType) +{ + constexpr int64_t A_UID = 10; + constexpr int64_t B_UID = 11; + constexpr int64_t C_UID = 12; + const std::vector dims = {4, 8, 2, 2}; + const std::vector strides = generateStrides(dims); + + auto graphBuilder = createMatmulGraph(A_UID, + B_UID, + C_UID, + dims, + strides, + dims, + strides, + dims, + strides, + aType, + bType, + cType, + computeType); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper graphWrapper( + graphBuilder.GetBufferPointer(), graphBuilder.GetSize()); + + return GpuMatmulSignatureKey(graphWrapper.getNode(0), graphWrapper.getTensorMap(), computeType); +} + +// Builds a batchnorm inference key from a real flatbuffer node and tensor map. +GpuPlanRegistrySignatureKey makeBatchnormInferenceKey(DataType inputType) +{ + auto builder = hipdnn_test_sdk::utilities::createValidBatchnormInferenceGraph( + {48, 16, 4, 1}, {2, 3, 4, 4}, inputType); + const hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper graphWrapper( + builder.GetBufferPointer(), builder.GetSize()); + + return GpuBatchnormFwdInfSignatureKey(graphWrapper.getNode(0), + graphWrapper.getTensorMap(), + graphWrapper.getNode(0).compute_data_type()); +} + +} // namespace + +TEST(TestGpuPlanRegistrySignatureKey, EqualKeysOfSameTypeAreEqualAndHashEqual) +{ + const auto key1 + = makeMatmulKey(DataType::FLOAT, DataType::HALF, DataType::FLOAT, DataType::FLOAT); + const auto key2 + = makeMatmulKey(DataType::FLOAT, DataType::HALF, DataType::FLOAT, DataType::FLOAT); + + const GpuPlanRegistrySignatureKeyHash hasher; + const GpuPlanRegistrySignatureKeyEqual eq; + + EXPECT_TRUE(eq(key1, key2)); + EXPECT_TRUE(eq(key2, key1)); + EXPECT_TRUE(eq(key1, key1)); + EXPECT_EQ(hasher(key1), hasher(key2)); + + const auto bnKey1 = makeBatchnormInferenceKey(DataType::FLOAT); + const auto bnKey2 = makeBatchnormInferenceKey(DataType::FLOAT); + + EXPECT_TRUE(eq(bnKey1, bnKey2)); + EXPECT_EQ(hasher(bnKey1), hasher(bnKey2)); +} + +TEST(TestGpuPlanRegistrySignatureKey, KeysOfSameTypeWithDifferentDataAreNotEqual) +{ + const auto base + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::FLOAT, DataType::FLOAT); + const auto diffA + = makeMatmulKey(DataType::HALF, DataType::FLOAT, DataType::FLOAT, DataType::FLOAT); + const auto diffB + = makeMatmulKey(DataType::FLOAT, DataType::HALF, DataType::FLOAT, DataType::FLOAT); + const auto diffC + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::HALF, DataType::FLOAT); + const auto diffCompute + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::FLOAT, DataType::HALF); + + const GpuPlanRegistrySignatureKeyEqual eq; + + EXPECT_FALSE(eq(base, diffA)); + EXPECT_FALSE(eq(base, diffB)); + EXPECT_FALSE(eq(base, diffC)); + EXPECT_FALSE(eq(base, diffCompute)); + EXPECT_FALSE(eq(diffA, base)); + + const auto bnFloat = makeBatchnormInferenceKey(DataType::FLOAT); + const auto bnHalf = makeBatchnormInferenceKey(DataType::HALF); + + EXPECT_FALSE(eq(bnFloat, bnHalf)); + EXPECT_FALSE(eq(bnHalf, bnFloat)); +} + +TEST(TestGpuPlanRegistrySignatureKey, KeysOfDifferentTypesAreNotEqual) +{ + const auto matmulKey + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::FLOAT, DataType::FLOAT); + const auto batchnormKey = makeBatchnormInferenceKey(DataType::FLOAT); + + const GpuPlanRegistrySignatureKeyEqual eq; + + ASSERT_NE(matmulKey.index(), batchnormKey.index()); + EXPECT_FALSE(eq(matmulKey, batchnormKey)); + EXPECT_FALSE(eq(batchnormKey, matmulKey)); +} + +TEST(TestGpuPlanRegistrySignatureKey, UnorderedMapLookupFindsOnlyMatchingKey) +{ + const auto matmulFloat + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::FLOAT, DataType::FLOAT); + const auto matmulHalf + = makeMatmulKey(DataType::HALF, DataType::HALF, DataType::HALF, DataType::FLOAT); + const auto bnFloat = makeBatchnormInferenceKey(DataType::FLOAT); + const auto bnHalf = makeBatchnormInferenceKey(DataType::HALF); + + std::unordered_map + registry; + registry.emplace(matmulFloat, 1); + registry.emplace(bnFloat, 2); + + // Re-built, equal keys find the inserted entries. + const auto matmulFloatAgain + = makeMatmulKey(DataType::FLOAT, DataType::FLOAT, DataType::FLOAT, DataType::FLOAT); + const auto bnFloatAgain = makeBatchnormInferenceKey(DataType::FLOAT); + + const auto matmulIt = registry.find(matmulFloatAgain); + ASSERT_NE(matmulIt, registry.end()); + EXPECT_EQ(matmulIt->second, 1); + + const auto bnIt = registry.find(bnFloatAgain); + ASSERT_NE(bnIt, registry.end()); + EXPECT_EQ(bnIt->second, 2); + + // Different-valued keys of the same type are not found. + EXPECT_EQ(registry.find(matmulHalf), registry.end()); + EXPECT_EQ(registry.find(bnHalf), registry.end()); + EXPECT_EQ(registry.size(), 2U); +} diff --git a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/cpu_graph_executor/detail/PlanRegistrySignatureKey.hpp b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/cpu_graph_executor/detail/PlanRegistrySignatureKey.hpp index fe1e1e48e2df..a1839df19734 100644 --- a/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/cpu_graph_executor/detail/PlanRegistrySignatureKey.hpp +++ b/projects/hipdnn/test_sdk/include/hipdnn_test_sdk/utilities/cpu_graph_executor/detail/PlanRegistrySignatureKey.hpp @@ -5,6 +5,7 @@ #include #include +#include #include #include @@ -77,22 +78,18 @@ struct PlanRegistrySignatureKeyHash struct PlanRegistrySignatureKeyEqual { - template - bool operator()([[maybe_unused]] const T& a, [[maybe_unused]] const U& b) const - { - return false; - } - - template - bool operator()(const T& a, const T& b) const - { - return a == b; - } - - // NOLINTNEXTLINE(readability-redundant-casting) + // Compare the active alternatives' indices first, then visit only `a` and fetch the matching + // alternative from `b`. Visiting both variants would instantiate one visitor call per pair of + // alternatives (N*N) instead of N. bool operator()(const PlanRegistrySignatureKey& a, const PlanRegistrySignatureKey& b) const { - return std::visit(*this, a, b); + if(a.index() != b.index()) + { + return false; + } + + return std::visit( + [&b](const auto& x) { return x == std::get>(b); }, a); } }; From 675adabe8f711aa021969fa4d33a7bce6a23366d Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 09:05:38 -0600 Subject: [PATCH 02/16] refactor(integration-tests): compile the CPU reference adapter and factory once --- .../integration-tests/CMakeLists.txt | 3 ++ .../CpuReferenceGraphExecutorAdapter.cpp | 52 +++++++++++++++++++ .../CpuReferenceGraphExecutorAdapter.hpp | 51 ++++++++---------- .../harness/ReferenceGraphExecutorFactory.cpp | 37 +++++++++++++ .../harness/ReferenceGraphExecutorFactory.hpp | 27 +--------- ...tegrationConvForwardSerializeRoundTrip.cpp | 2 + .../moe/IntegrationGpuMoeGroupedMatmul.cpp | 1 + .../moe/IntegrationGpuMoeGroupedMatmulBwd.cpp | 1 + .../IntegrationGpuResampleBackward.cpp | 2 + .../sdpa/IntegrationGpuSdpaFwdInference.cpp | 1 + .../integration-tests/tests/CMakeLists.txt | 16 +++--- .../tests/TestErrorPaths.cpp | 34 ++++++++++-- 12 files changed, 164 insertions(+), 63 deletions(-) create mode 100644 dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.cpp create mode 100644 dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.cpp diff --git a/dnn-providers/integration-tests/CMakeLists.txt b/dnn-providers/integration-tests/CMakeLists.txt index ff1d43a5cf3f..5d9b1deaa7a2 100644 --- a/dnn-providers/integration-tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/CMakeLists.txt @@ -223,7 +223,9 @@ set(INTEGRATION_TESTS_EXE hipdnn_integration_tests) add_executable(${INTEGRATION_TESTS_EXE} src/main.cpp + src/harness/CpuReferenceGraphExecutorAdapter.cpp src/harness/ReferenceExecutorPool.cpp + src/harness/ReferenceGraphExecutorFactory.cpp src/harness/bundle/FrontendGraphEngineRunner.cpp src/harness/bundle/GraphSession.cpp src/harness/bundle/HarnessDependencies.cpp @@ -255,6 +257,7 @@ set(GOLDEN_DATA_TESTS_EXE hipdnn_golden_data_tests) add_executable(${GOLDEN_DATA_TESTS_EXE} src/golden_data_main.cpp + src/harness/CpuReferenceGraphExecutorAdapter.cpp src/harness/ReferenceExecutorPool.cpp src/harness/bundle/GraphSession.cpp src/harness/bundle/OutputComparison.cpp diff --git a/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.cpp b/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.cpp new file mode 100644 index 000000000000..405c2524cb57 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.cpp @@ -0,0 +1,52 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" + +#include + +#include +#include + +namespace hipdnn_integration_tests +{ + +struct CpuReferenceGraphExecutorAdapter::Impl +{ + hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor executor; +}; + +CpuReferenceGraphExecutorAdapter::CpuReferenceGraphExecutorAdapter() + : _impl(std::make_unique()) +{ +} + +CpuReferenceGraphExecutorAdapter::~CpuReferenceGraphExecutorAdapter() = default; + +bool CpuReferenceGraphExecutorAdapter::isApplicable(void* graphBuffer, size_t size) +{ + return _impl->executor.isApplicable(graphBuffer, size); +} + +void CpuReferenceGraphExecutorAdapter::execute( + void* graphBuffer, size_t size, const std::unordered_map& variantPack) +{ + // The test_sdk CPU executor throws CpuReferenceNotApplicableError for + // capability misses (no plan for this node type, or plan builder says + // not applicable). We translate that into a ReferenceCapabilityError + // (case A) so the harness can skip. Any other exception is a genuine + // runtime failure (case C) and propagates unchanged — the harness + // reports it as RUNTIME_ERROR → FAIL(). + try + { + _impl->executor.execute(graphBuffer, size, variantPack); + } + catch(const hipdnn_test_sdk::utilities::CpuReferenceNotApplicableError& e) + { + throw ReferenceCapabilityError(std::string("CPU reference executor could not run " + "this graph: ") + + e.what()); + } +} + +} // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.hpp b/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.hpp index d5260424c741..b888cadb8297 100644 --- a/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.hpp +++ b/dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.hpp @@ -3,11 +3,10 @@ #pragma once -#include -#include - -#include -#include +#include +#include +#include +#include #include "IReferenceGraphExecutor.hpp" #include "ReferenceCapabilityError.hpp" @@ -15,35 +14,28 @@ namespace hipdnn_integration_tests { +// Adapts the test_sdk CPU reference executor to IReferenceGraphExecutor. The +// executor is held behind a pimpl so the heavy test_sdk executor header is +// compiled once (CpuReferenceGraphExecutorAdapter.cpp) instead of in every +// translation unit that needs a CPU reference. class CpuReferenceGraphExecutorAdapter : public IReferenceGraphExecutor { public: - bool isApplicable(void* graphBuffer, size_t size) override - { - return _executor.isApplicable(graphBuffer, size); - } + CpuReferenceGraphExecutorAdapter(); + ~CpuReferenceGraphExecutorAdapter() override; + + CpuReferenceGraphExecutorAdapter(const CpuReferenceGraphExecutorAdapter&) = delete; + CpuReferenceGraphExecutorAdapter& operator=(const CpuReferenceGraphExecutorAdapter&) = delete; + CpuReferenceGraphExecutorAdapter(CpuReferenceGraphExecutorAdapter&&) = delete; + CpuReferenceGraphExecutorAdapter& operator=(CpuReferenceGraphExecutorAdapter&&) = delete; + bool isApplicable(void* graphBuffer, size_t size) override; + + // Translates the test_sdk CpuReferenceNotApplicableError into a + // ReferenceCapabilityError; any other exception propagates unchanged. void execute(void* graphBuffer, size_t size, - const std::unordered_map& variantPack) override - { - // The test_sdk CPU executor throws CpuReferenceNotApplicableError for - // capability misses (no plan for this node type, or plan builder says - // not applicable). We translate that into a ReferenceCapabilityError - // (case A) so the harness can skip. Any other exception is a genuine - // runtime failure (case C) and propagates unchanged — the harness - // reports it as RUNTIME_ERROR → FAIL(). - try - { - _executor.execute(graphBuffer, size, variantPack); - } - catch(const hipdnn_test_sdk::utilities::CpuReferenceNotApplicableError& e) - { - throw ReferenceCapabilityError(std::string("CPU reference executor could not run " - "this graph: ") - + e.what()); - } - } + const std::unordered_map& variantPack) override; bool requiresDeviceMemory() const override { @@ -51,7 +43,8 @@ class CpuReferenceGraphExecutorAdapter : public IReferenceGraphExecutor } private: - hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor _executor; + struct Impl; + std::unique_ptr _impl; }; } // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.cpp b/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.cpp new file mode 100644 index 000000000000..236fbca311fa --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.cpp @@ -0,0 +1,37 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "harness/ReferenceGraphExecutorFactory.hpp" + +#include + +#include + +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" +#include "harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp" + +namespace hipdnn_integration_tests +{ + +std::unique_ptr + ReferenceGraphExecutorFactory::create(ReferenceExecutorType type) +{ + switch(type) + { + case ReferenceExecutorType::CPU: + HIPDNN_PLUGIN_LOG_TRACE("ReferenceGraphExecutorFactory: creating CPU reference executor"); + return std::make_unique(); + case ReferenceExecutorType::GPU: + HIPDNN_PLUGIN_LOG_TRACE("ReferenceGraphExecutorFactory: creating GPU reference executor"); + return std::make_unique(); + default: + throw std::runtime_error("Unknown reference executor type"); + } +} + +std::unique_ptr ReferenceGraphExecutorFactory::createFromConfig() +{ + return create(TestConfig::get().getReferenceExecutorType()); +} + +} // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.hpp b/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.hpp index b7f451c49f64..02aa5135d6c0 100644 --- a/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.hpp +++ b/dnn-providers/integration-tests/src/harness/ReferenceGraphExecutorFactory.hpp @@ -4,14 +4,9 @@ #pragma once #include -#include -#include - -#include "CpuReferenceGraphExecutorAdapter.hpp" #include "IReferenceGraphExecutor.hpp" #include "TestConfig.hpp" -#include "gpu-graph-executor/GpuReferenceGraphExecutor.hpp" namespace hipdnn_integration_tests { @@ -19,28 +14,10 @@ namespace hipdnn_integration_tests class ReferenceGraphExecutorFactory { public: - static std::unique_ptr create(ReferenceExecutorType type) - { - switch(type) - { - case ReferenceExecutorType::CPU: - HIPDNN_PLUGIN_LOG_TRACE( - "ReferenceGraphExecutorFactory: creating CPU reference executor"); - return std::make_unique(); - case ReferenceExecutorType::GPU: - HIPDNN_PLUGIN_LOG_TRACE( - "ReferenceGraphExecutorFactory: creating GPU reference executor"); - return std::make_unique(); - default: - throw std::runtime_error("Unknown reference executor type"); - } - } + static std::unique_ptr create(ReferenceExecutorType type); /// Convenience: create from TestConfig - static std::unique_ptr createFromConfig() - { - return create(TestConfig::get().getReferenceExecutorType()); - } + static std::unique_ptr createFromConfig(); }; } // namespace hipdnn_integration_tests diff --git a/dnn-providers/integration-tests/src/integration-tests/conv/IntegrationConvForwardSerializeRoundTrip.cpp b/dnn-providers/integration-tests/src/integration-tests/conv/IntegrationConvForwardSerializeRoundTrip.cpp index 2134be9df7c6..5efb5d9f3d88 100644 --- a/dnn-providers/integration-tests/src/integration-tests/conv/IntegrationConvForwardSerializeRoundTrip.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/conv/IntegrationConvForwardSerializeRoundTrip.cpp @@ -12,6 +12,8 @@ #include +#include + #include #include #include diff --git a/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmul.cpp b/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmul.cpp index 12375e587509..766f065a8b0d 100644 --- a/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmul.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmul.cpp @@ -5,6 +5,7 @@ #include #include +#include #include #include diff --git a/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmulBwd.cpp b/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmulBwd.cpp index 12cd7b3293be..7bcf0d6b9882 100644 --- a/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmulBwd.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/moe/IntegrationGpuMoeGroupedMatmulBwd.cpp @@ -11,6 +11,7 @@ #include #include +#include #include #include diff --git a/dnn-providers/integration-tests/src/integration-tests/resample/IntegrationGpuResampleBackward.cpp b/dnn-providers/integration-tests/src/integration-tests/resample/IntegrationGpuResampleBackward.cpp index 77d99fd74f46..7e3bea300e93 100644 --- a/dnn-providers/integration-tests/src/integration-tests/resample/IntegrationGpuResampleBackward.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/resample/IntegrationGpuResampleBackward.cpp @@ -1,6 +1,8 @@ // Copyright © Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT +#include + #include #include #include diff --git a/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationGpuSdpaFwdInference.cpp b/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationGpuSdpaFwdInference.cpp index b2e34bcbd033..f8ebc0cfbdf1 100644 --- a/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationGpuSdpaFwdInference.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationGpuSdpaFwdInference.cpp @@ -4,6 +4,7 @@ #ifdef HIPDNN_ENABLE_SDPA #include +#include #include #include diff --git a/dnn-providers/integration-tests/tests/CMakeLists.txt b/dnn-providers/integration-tests/tests/CMakeLists.txt index dbadb3b5987b..2fc735bc4cf5 100644 --- a/dnn-providers/integration-tests/tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/tests/CMakeLists.txt @@ -3,14 +3,18 @@ # NOTE: FrontendGraphEngineRunner.cpp, ReferenceExecutorPool.cpp and # HarnessDependencies.cpp are deliberately absent. They are the only places that -# touch the frontend Graph, a hipdnnHandle_t, or a real reference executor, so -# leaving them out makes "these unit tests never reach an engine or a device" a -# link-time guarantee rather than a convention. GoldenDataRegistration.cpp is absent -# for the same reason: it registers gtest tests and reaches the shared reference -# executors. Its decisions live in GoldenDataPlan.cpp and GoldenOutputProbe.cpp, -# which are linked and tested here. +# touch the frontend Graph or a hipdnnHandle_t, or hand the harness a pooled reference +# executor. CpuReferenceGraphExecutorAdapter.cpp and ReferenceGraphExecutorFactory.cpp +# ARE linked: the direct executor tests construct reference executors themselves. +# Leaving the pool, HarnessDependencies.cpp and GoldenDataRegistration.cpp out keeps +# "harness code paths cannot get a pooled executor" a link-time guarantee rather than +# a convention. GoldenDataRegistration.cpp is absent for the same reason: it registers +# gtest tests and reaches the shared reference executors. Its decisions live in +# GoldenDataPlan.cpp and GoldenOutputProbe.cpp, which are linked and tested here. add_executable(hipdnn_integration_tests_unit_tests main.cpp + ../src/harness/CpuReferenceGraphExecutorAdapter.cpp + ../src/harness/ReferenceGraphExecutorFactory.cpp ../src/harness/bundle/GraphSession.cpp ../src/harness/bundle/IntegrationBundleVerificationHarness.cpp ../src/harness/bundle/OutputComparison.cpp diff --git a/dnn-providers/integration-tests/tests/TestErrorPaths.cpp b/dnn-providers/integration-tests/tests/TestErrorPaths.cpp index 99d40bd8a5e0..dad89b6bd672 100644 --- a/dnn-providers/integration-tests/tests/TestErrorPaths.cpp +++ b/dnn-providers/integration-tests/tests/TestErrorPaths.cpp @@ -8,16 +8,19 @@ #include #include +#include #include #include #include #include #include +#include #include #include #include +#include #include #include @@ -229,10 +232,12 @@ TEST_F(TestErrorPaths, RefCrashFails) EXPECT_THAT(refErrors.front(), ::testing::HasSubstr("stub: ref crashed on supported op")); } -TEST_F(TestErrorPaths, AdapterTranslatesNotApplicableToCapabilityError) +// Pins the exception type relationships the harness routing relies on: the CPU +// executor's CpuReferenceNotApplicableError is a plain std::runtime_error, NOT a +// ReferenceCapabilityError, so only the adapter's translation turns it into a skip. +// This test does not call the adapter; see CpuAdapterTranslatesNotApplicableToCapabilityError. +TEST_F(TestErrorPaths, NotApplicableAndCapabilityErrorTypeRelationships) { - const CpuReferenceGraphExecutorAdapter adapter; - hipdnn_test_sdk::utilities::CpuReferenceNotApplicableError notApplicable("stub"); EXPECT_TRUE(dynamic_cast(¬Applicable) != nullptr) << "CpuReferenceNotApplicableError must derive from std::runtime_error"; @@ -255,6 +260,29 @@ TEST_F(TestErrorPaths, AdapterTranslatesNotApplicableToCapabilityError) } } +// A CustomOp node has no CPU signature key, so the test_sdk executor throws +// CpuReferenceNotApplicableError from execute(); the adapter must surface that as a +// ReferenceCapabilityError so the harness skips instead of failing. +TEST_F(TestErrorPaths, CpuAdapterTranslatesNotApplicableToCapabilityError) +{ + auto builder = hipdnn_test_sdk::utilities::createValidCustomOpGraph(); + const std::unordered_map variantPack; + CpuReferenceGraphExecutorAdapter adapter; + + EXPECT_FALSE(adapter.isApplicable(builder.GetBufferPointer(), builder.GetSize())); + + try + { + adapter.execute(builder.GetBufferPointer(), builder.GetSize(), variantPack); + FAIL() << "execute() must throw for a graph the CPU executor cannot run"; + } + catch(const ReferenceCapabilityError& e) + { + EXPECT_THAT(std::string(e.what()), + ::testing::StartsWith("CPU reference executor could not run")); + } +} + TEST_F(TestErrorPaths, GenericRuntimeErrorNotCaughtAsNotApplicable) { try From 26c18c4043a248401abef7ed56a57311e20d35cc Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 09:08:51 -0600 Subject: [PATCH 03/16] refactor(integration-tests): use the CPU reference adapter in plan tests --- .../src/harness/IReferenceExecutors.hpp | 9 +++++---- .../src/harness/ReferenceCapabilityError.hpp | 7 ++++--- .../sdpa/IntegrationCpuSdpaFwdBwdChained.cpp | 7 ++++--- .../tests/TestGpuBatchnormFwdInfPlan.cpp | 4 ++-- .../tests/TestGpuBatchnormFwdInfVariancePlan.cpp | 4 ++-- .../tests/TestGpuBatchnormFwdTrainPlan.cpp | 4 ++-- .../tests/TestGpuConvolutionFwdPlan.cpp | 4 ++-- .../tests/TestGpuLayernormBwdPlan.cpp | 4 ++-- .../tests/TestGpuLayernormFwdPlan.cpp | 4 ++-- .../integration-tests/tests/TestGpuMatmulPlan.cpp | 4 ++-- .../integration-tests/tests/TestGpuRMSNormPlan.cpp | 6 +++--- .../tests/TestGpuReferenceGraphExecutor.cpp | 9 ++++++--- .../integration-tests/tests/TestGpuResamplePlan.cpp | 7 ++++--- .../tests/TestVerificationPaths.cpp | 12 ++++++------ 14 files changed, 46 insertions(+), 39 deletions(-) diff --git a/dnn-providers/integration-tests/src/harness/IReferenceExecutors.hpp b/dnn-providers/integration-tests/src/harness/IReferenceExecutors.hpp index e625e7a56c7f..fcab015f5a8d 100644 --- a/dnn-providers/integration-tests/src/harness/IReferenceExecutors.hpp +++ b/dnn-providers/integration-tests/src/harness/IReferenceExecutors.hpp @@ -11,10 +11,11 @@ namespace hipdnn_integration_tests /// Owns the run's reference executors and lends them out. /// -/// A container rather than a factory on purpose. The executors are reusable — a -/// GpuReferenceGraphExecutor's plan-builder registry is a lookup table with no -/// per-graph state — so building a fresh one per bundle threw the table away -/// hundreds of times per run and made the caller own something it only borrows. +/// A container rather than a factory on purpose. The executors are reusable — the +/// plan-builder registry inside a GpuReferenceGraphExecutor or a +/// CpuReferenceGraphExecutorAdapter is a lookup table with no per-graph state — so +/// building a fresh one per bundle threw the table away hundreds of times per run +/// and made the caller own something it only borrows. /// /// Lifetime: the returned reference stays valid for the life of the container, so /// the production container is held for the whole process. diff --git a/dnn-providers/integration-tests/src/harness/ReferenceCapabilityError.hpp b/dnn-providers/integration-tests/src/harness/ReferenceCapabilityError.hpp index 498c70f03700..fd2073b00ad1 100644 --- a/dnn-providers/integration-tests/src/harness/ReferenceCapabilityError.hpp +++ b/dnn-providers/integration-tests/src/harness/ReferenceCapabilityError.hpp @@ -29,9 +29,10 @@ namespace hipdnn_integration_tests // // NOTE: the GPU reference executor (our code) throws this directly at its // capability-miss sites. The CPU reference executor (test_sdk) throws -// CpuReferenceNotApplicableError for capability misses; the -// CpuReferenceGraphExecutorAdapter translates that into a ReferenceCapabilityError. -// Other test_sdk exceptions propagate as genuine runtime failures (case C). +// CpuReferenceNotApplicableError for capability misses; CpuReferenceGraphExecutorAdapter +// (CpuReferenceGraphExecutorAdapter.cpp) catches that and rethrows a +// ReferenceCapabilityError, so code behind IReferenceGraphExecutor only ever sees this +// type. Other test_sdk exceptions propagate as genuine runtime failures (case C). class ReferenceCapabilityError : public std::runtime_error { public: diff --git a/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationCpuSdpaFwdBwdChained.cpp b/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationCpuSdpaFwdBwdChained.cpp index 18a9cc101f4f..2b78e5a0c13c 100644 --- a/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationCpuSdpaFwdBwdChained.cpp +++ b/dnn-providers/integration-tests/src/integration-tests/sdpa/IntegrationCpuSdpaFwdBwdChained.cpp @@ -9,9 +9,10 @@ #include #include #include -#include #include +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" + namespace { @@ -125,7 +126,7 @@ TEST_F(IntegrationCpuSdpaFwdBwdChained, Bf16NoMaskProducesFiniteGradients) // Execute forward CPU reference. auto fwdVariantPack = fwdBundle.toHostVariantPack(); - hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute( + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter().execute( fwdBuilder.GetBufferPointer(), fwdBuilder.GetSize(), fwdVariantPack); // Validate forward outputs are finite. @@ -174,7 +175,7 @@ TEST_F(IntegrationCpuSdpaFwdBwdChained, Bf16NoMaskProducesFiniteGradients) // Execute backward CPU reference. auto bwdVariantPack = bwdBundle.toHostVariantPack(); - hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute( + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter().execute( bwdBuilder.GetBufferPointer(), bwdBuilder.GetSize(), bwdVariantPack); // Validate backward outputs are finite and non-zero. diff --git a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfPlan.cpp index 1d772fe24669..c37992faa5fb 100644 --- a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfPlan.cpp @@ -6,6 +6,7 @@ #include #include +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuBatchnormFwdInfPlan.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" #include @@ -13,7 +14,6 @@ #include #include #include -#include using namespace hipdnn_flatbuffers_sdk::data_objects; using namespace hipdnn_integration_tests::gpu_graph_executor::detail; @@ -269,7 +269,7 @@ void runPlanExecuteVsCpuRef(const std::vector& dims, cpuVariantPack[nodeAttributes->mean_tensor_uid()] = meanTensor.rawHostData(); cpuVariantPack[nodeAttributes->inv_variance_tensor_uid()] = invVarianceTensor.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); cpuY.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfVariancePlan.cpp b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfVariancePlan.cpp index e83a8b463c39..a3ee1d02ed97 100644 --- a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfVariancePlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdInfVariancePlan.cpp @@ -4,6 +4,7 @@ #include #include "ScalarTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" #include #include @@ -12,7 +13,6 @@ #include #include #include -#include #include #include "harness/gpu-graph-executor/detail/GpuBatchnormFwdInfVariancePlan.hpp" @@ -391,7 +391,7 @@ void runPlanExecuteVsCpuRef(const std::vector& dims, cpuVariantPack[nodeAttributes->variance_tensor_uid()] = varianceTensor.rawHostData(); cpuVariantPack[nodeAttributes->epsilon_tensor_uid()] = epsilonTensor.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); cpuY.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdTrainPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdTrainPlan.cpp index d2f71a25782c..065e5998300c 100644 --- a/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdTrainPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuBatchnormFwdTrainPlan.cpp @@ -7,13 +7,13 @@ #include #include "BatchnormFwdTrainGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuBatchnormFwdTrainPlan.hpp" #include #include #include #include #include -#include using namespace hipdnn_flatbuffers_sdk::data_objects; using namespace hipdnn_integration_tests::gpu_graph_executor::detail; @@ -513,7 +513,7 @@ void runPlanExecuteVsCpuRef(const std::vector& dims, const TensorLayout cpuVariantPack[nodeAttributes->next_running_variance_tensor_uid().value()] = cpuNextRunningVariance.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(builder.GetBufferPointer(), builder.GetSize(), cpuVariantPack); cpuOutput.markHostModified(); cpuMean.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuConvolutionFwdPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuConvolutionFwdPlan.cpp index 66a37d64e405..3d522c9fa04d 100644 --- a/dnn-providers/integration-tests/tests/TestGpuConvolutionFwdPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuConvolutionFwdPlan.cpp @@ -16,10 +16,10 @@ #include #include #include -#include #include "ConvolutionFwdGraphTestUtils.hpp" #include "RaggedGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuConvolutionFwdPlan.hpp" #include "harness/gpu-graph-executor/detail/GpuConvolutionFwdSignatureKey.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" @@ -275,7 +275,7 @@ void runPlanExecuteVsCpuRef(const std::vector& xDims, cpuVariantPack[W_UID] = cpuW.rawHostData(); cpuVariantPack[Y_UID] = cpuY.rawHostData(); - hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); // Compare diff --git a/dnn-providers/integration-tests/tests/TestGpuLayernormBwdPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuLayernormBwdPlan.cpp index 08430bdc4ca8..e0c248e40b63 100644 --- a/dnn-providers/integration-tests/tests/TestGpuLayernormBwdPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuLayernormBwdPlan.cpp @@ -16,10 +16,10 @@ #include #include #include -#include #include "LayernormBwdGraphTestUtils.hpp" #include "RaggedGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuLayernormBwdPlan.hpp" #include "harness/gpu-graph-executor/detail/GpuLayernormBwdSignatureKey.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" @@ -455,7 +455,7 @@ void runPlanExecuteVsCpuRef(const std::vector& ioDims, cpuVariantPack[INV_VARIANCE_UID] = rstdTensor.rawHostData(); } - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); cpuDx.markHostModified(); cpuDscale.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuLayernormFwdPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuLayernormFwdPlan.cpp index 17ef8f35e9be..58ffbc792619 100644 --- a/dnn-providers/integration-tests/tests/TestGpuLayernormFwdPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuLayernormFwdPlan.cpp @@ -16,10 +16,10 @@ #include #include #include -#include #include "LayernormFwdGraphTestUtils.hpp" #include "RaggedGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuLayernormFwdPlan.hpp" #include "harness/gpu-graph-executor/detail/GpuLayernormFwdSignatureKey.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" @@ -404,7 +404,7 @@ void runPlanExecuteVsCpuRef(const std::vector& ioDims, cpuVariantPack[MEAN_UID] = cpuMean.rawHostData(); cpuVariantPack[INV_VARIANCE_UID] = cpuRstd.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); cpuY.markHostModified(); cpuMean.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuMatmulPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuMatmulPlan.cpp index 18da23687995..7338c7bb3180 100644 --- a/dnn-providers/integration-tests/tests/TestGpuMatmulPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuMatmulPlan.cpp @@ -10,11 +10,11 @@ #include #include #include -#include #include #include "MatmulGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuMatmulPlan.hpp" #include "harness/gpu-graph-executor/detail/GpuMatmulSignatureKey.hpp" #include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" @@ -522,7 +522,7 @@ void runPlanExecuteVsCpuRef(const std::vector& aDims, cpuVariantPack[B_UID] = bTensor.rawHostData(); cpuVariantPack[C_UID] = cpuC.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); cpuC.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuRMSNormPlan.cpp b/dnn-providers/integration-tests/tests/TestGpuRMSNormPlan.cpp index 6d5e57289217..cef7cb02d060 100644 --- a/dnn-providers/integration-tests/tests/TestGpuRMSNormPlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuRMSNormPlan.cpp @@ -9,13 +9,13 @@ #include "RaggedGraphTestUtils.hpp" #include "ScalarTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuRMSNormPlan.hpp" #include #include #include #include #include -#include using namespace hipdnn_flatbuffers_sdk::data_objects; using namespace hipdnn_integration_tests::gpu_graph_executor::detail; @@ -482,7 +482,7 @@ void runFwdPlanExecuteVsCpuRef(const std::vector& ioDims, cpuVariantPack[nodeAttributes->scale_tensor_uid()] = scaleTensor.rawHostData(); cpuVariantPack[nodeAttributes->y_tensor_uid()] = cpuY.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(builder.GetBufferPointer(), builder.GetSize(), cpuVariantPack); cpuY.markHostModified(); @@ -657,7 +657,7 @@ void runBwdPlanExecuteVsCpuRef(const std::vector& ioDims, cpuVariantPack[nodeAttributes->dscale_tensor_uid()] = dCpuScaleTensor.rawHostData(); cpuVariantPack[nodeAttributes->dbias_tensor_uid().value()] = dCpuBiasTensor.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(builder.GetBufferPointer(), builder.GetSize(), cpuVariantPack); dxCpuTensor.markHostModified(); dCpuScaleTensor.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp b/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp index 2fdb4f2cfa1f..5502f88a8cd2 100644 --- a/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp @@ -7,8 +7,11 @@ #include #include #include +#include #include #include +#include +#include #include #include #include @@ -19,13 +22,13 @@ #include #include #include -#include #include "ConvolutionFwdGraphTestUtils.hpp" #include "LayernormBwdGraphTestUtils.hpp" #include "LayernormFwdGraphTestUtils.hpp" #include "MatmulGraphTestUtils.hpp" #include "RaggedGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/ReferenceCapabilityError.hpp" #include "harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp" @@ -35,8 +38,8 @@ namespace using namespace hipdnn_flatbuffers_sdk::data_objects; using namespace hipdnn_integration_tests::test_utils; using namespace hipdnn_gpu_ref::common::gpu_fp_reference_tensor; +using hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter; using hipdnn_integration_tests::gpu_graph_executor::GpuReferenceGraphExecutor; -using hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor; // Creates a minimal pointwise graph with two FLOAT tensors (input + output). // The pointwise operation is RELU_FWD but the dummy plan ignores the operation. @@ -352,7 +355,7 @@ void runConvFwdExecutorVsCpu(const std::vector& xDims, cpuVariantPack[W_UID] = cpuW.rawHostData(); cpuVariantPack[Y_UID] = cpuY.rawHostData(); - CpuReferenceGraphExecutor cpuExecutor; + CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(graphBuilder.GetBufferPointer(), graphBuilder.GetSize(), cpuVariantPack); // Compare GPU executor output against CPU executor output diff --git a/dnn-providers/integration-tests/tests/TestGpuResamplePlan.cpp b/dnn-providers/integration-tests/tests/TestGpuResamplePlan.cpp index 0276ca8ba908..dfdf92976b5e 100644 --- a/dnn-providers/integration-tests/tests/TestGpuResamplePlan.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuResamplePlan.cpp @@ -10,13 +10,14 @@ #include #include "RaggedGraphTestUtils.hpp" +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/gpu-graph-executor/detail/GpuResamplePlan.hpp" #include #include +#include #include #include #include -#include using namespace hipdnn_flatbuffers_sdk::data_objects; using namespace hipdnn_integration_tests::gpu_graph_executor::detail; @@ -556,7 +557,7 @@ void runResampleFwdPlanExecuteVsCpuRef(const std::vector& xDims, cpuVariantPack[INDEX_UID] = cpuIndex.rawHostData(); } - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(builder.GetBufferPointer(), builder.GetSize(), cpuVariantPack); cpuY.markHostModified(); cpuIndex.markHostModified(); @@ -716,7 +717,7 @@ void runResampleBwdPlanExecuteVsCpuRef(const std::vector& dxDims, cpuVariantPack[INDEX_UID] = indexTensor.rawHostData(); } - CpuReferenceGraphExecutor cpuExecutor; + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter cpuExecutor; cpuExecutor.execute(builder.GetBufferPointer(), builder.GetSize(), cpuVariantPack); cpuDx.markHostModified(); diff --git a/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp b/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp index 1b8dfc3dc7de..05f334402b11 100644 --- a/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp +++ b/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp @@ -12,8 +12,8 @@ #include #include #include -#include +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/bundle/BundleDiscovery.hpp" #include "harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp" @@ -91,10 +91,10 @@ void writeMinimalBatchNormBundle(const std::filesystem::path& dir, const std::st // --------------------------------------------------------------------------- // Path 1 — CPU Reference against golden data (the "LayerNorm default" path) // -// For ops without bundles AND without a GPU reference executor, the -// existing TestCpuFpReference* tests use CpuReferenceGraphExecutor as the -// ground truth. This test proves the golden comparison pipeline works -// end-to-end by running the same CPU executor against a real batch-norm +// For ops without bundles AND without a GPU reference executor, the CPU reference +// is the ground truth. This test proves the golden comparison pipeline works +// end-to-end by running the CPU reference through CpuReferenceGraphExecutorAdapter +// (the same CPU path the verification harness uses) against a real batch-norm // bundle and comparing the result to the golden tensor data on disk. // --------------------------------------------------------------------------- TEST(TestGoldenVerificationCpuRefFp32, BatchNormSmallMatchesGoldenData) @@ -117,7 +117,7 @@ TEST(TestGoldenVerificationCpuRefFp32, BatchNormSmallMatchesGoldenData) auto goldenOutputs = graphAndTensors.extractAndClearOutputTensorData(); auto hostBuffers = graphAndTensors.hostBufferMap(); - hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute( + hipdnn_integration_tests::CpuReferenceGraphExecutorAdapter().execute( graphAndTensors.graphBuffer.data(), graphAndTensors.graphBuffer.size(), hostBuffers); verifyGoldenComparison(graphAndTensors, goldenOutputs, 1e-5f); From b07d928ea7cb8d0de321bb55508a2877acd91fcf Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 09:13:47 -0600 Subject: [PATCH 04/16] refactor(integration-tests): compile the GPU reference executor once --- .../integration-tests/CMakeLists.txt | 2 + .../GpuReferenceGraphExecutor.cpp | 222 ++++++++++++++++++ .../GpuReferenceGraphExecutor.hpp | 206 ++-------------- .../integration-tests/tests/CMakeLists.txt | 6 +- .../tests/TestGpuReferenceGraphExecutor.cpp | 1 + .../tests/TestVerificationPaths.cpp | 1 + 6 files changed, 248 insertions(+), 190 deletions(-) create mode 100644 dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp diff --git a/dnn-providers/integration-tests/CMakeLists.txt b/dnn-providers/integration-tests/CMakeLists.txt index 5d9b1deaa7a2..047af95bb4bd 100644 --- a/dnn-providers/integration-tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/CMakeLists.txt @@ -238,6 +238,7 @@ add_executable(${INTEGRATION_TESTS_EXE} src/harness/bundle/SupportVerdict.cpp src/harness/bundle/SweepManifestCache.cpp src/harness/bundle/VariantPackBuilder.cpp + src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp src/harness/input-init/FillInputs.cpp src/harness/input-init/SdpaBackwardInputs.cpp ) @@ -266,6 +267,7 @@ add_executable(${GOLDEN_DATA_TESTS_EXE} src/harness/bundle/SupportVerdict.cpp src/harness/bundle/SweepManifestCache.cpp src/harness/bundle/VariantPackBuilder.cpp + src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp src/harness/input-init/FillInputs.cpp src/harness/input-init/SdpaBackwardInputs.cpp src/harness/reference-validation/BundleReferenceValidationHarness.cpp diff --git a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp new file mode 100644 index 000000000000..082447ae6379 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp @@ -0,0 +1,222 @@ +// Copyright © Advanced Micro Devices, Inc., or its affiliates. +// SPDX-License-Identifier: MIT + +#include "harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp" + +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include "harness/ReferenceCapabilityError.hpp" +#include "harness/gpu-graph-executor/detail/GpuPlanBuilderRegistry.hpp" + +namespace hipdnn_integration_tests::gpu_graph_executor +{ + +namespace +{ + +std::unordered_map populateVariantPackWithMissingVirtualTensors( + const std::unordered_map& variantPack, + const std::unordered_map& + tensorMap, + std::vector>& virtualTensors) +{ + auto updatedVariantPack = variantPack; + + for(const auto& [id, attr] : tensorMap) + { + if(attr->virtual_() && updatedVariantPack.find(id) == updatedVariantPack.end()) + { + auto tensor = hipdnn_test_sdk::detail::createTensorFromAttribute(*attr); + virtualTensors.push_back(std::move(tensor)); + + updatedVariantPack[id] = virtualTensors.back()->rawDeviceData(); + } + } + return updatedVariantPack; +} + +detail::GpuPlanRegistrySignatureKey buildSignatureKey( + const hipdnn_flatbuffers_sdk::data_objects::Node& node, + const std::unordered_map& + tensorMap) +{ + using NodeAttrs = hipdnn_flatbuffers_sdk::data_objects::NodeAttributes; + + switch(node.attributes_type()) + { + case NodeAttrs::BatchnormAttributes: + return detail::GpuBatchnormFwdTrainSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::ConvolutionFwdAttributes: + return detail::GpuConvolutionFwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::LayernormAttributes: + return detail::GpuLayernormFwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::LayernormBackwardAttributes: + return detail::GpuLayernormBwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::MatmulAttributes: + return detail::GpuMatmulSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::PointwiseAttributes: + return detail::GpuPointwiseSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::ReductionAttributes: + return detail::GpuReductionSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::ResampleFwdAttributes: + return detail::GpuResampleFwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::ResampleBwdAttributes: + return detail::GpuResampleBwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::RMSNormAttributes: + return detail::GpuRMSNormFwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::RMSNormBackwardAttributes: + return detail::GpuRMSNormBwdSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::BatchnormInferenceAttributes: + return detail::GpuBatchnormFwdInfSignatureKey(node, tensorMap, node.compute_data_type()); + case NodeAttrs::BatchnormInferenceAttributesVarianceExt: + return detail::GpuBatchnormFwdInfVarianceSignatureKey( + node, tensorMap, node.compute_data_type()); + case NodeAttrs::SdpaAttributes: + { + // A ragged SDPA node (RFC-0014 packed layout) has a ragged_offset on a primary. Any + // ragged primary picks the ragged bucket, so a partly ragged node is rejected there + // instead of being run as dense. + const auto* sdpaAttributes = node.attributes_as_SdpaAttributes(); + if(sdpaAttributes != nullptr) + { + for(const auto uid : {sdpaAttributes->q_tensor_uid(), + sdpaAttributes->k_tensor_uid(), + sdpaAttributes->v_tensor_uid(), + sdpaAttributes->o_tensor_uid()}) + { + const auto it = tensorMap.find(uid); + if(it != tensorMap.end() && it->second != nullptr + && it->second->ragged_offset_tensor_uid().has_value()) + { + return detail::GpuSdpaRaggedFwdSignatureKey(node, tensorMap); + } + } + } + return detail::GpuSdpaFwdSignatureKey(node, tensorMap); + } + + // Node types with no GPU plan yet - throw descriptive error + case NodeAttrs::BatchnormBackwardAttributes: + case NodeAttrs::ConvolutionBwdAttributes: + case NodeAttrs::ConvolutionWrwAttributes: + case NodeAttrs::SdpaBackwardAttributes: + case NodeAttrs::BlockScaleDequantizeAttributes: + case NodeAttrs::BlockScaleQuantizeAttributes: + { + const std::string nodeName = node.name() == nullptr ? "unknown" : node.name()->str(); + throw ReferenceCapabilityError("GPU plan not yet implemented for node '" + nodeName + + "'. Register a GPU plan for this operation type."); + } + + case NodeAttrs::CustomOpAttributes: + throw ReferenceCapabilityError("GPU reference executor does not support custom operations"); + + default: + throw ReferenceCapabilityError("Unsupported node type for GPU signature key generation"); + } +} + +} // namespace + +struct GpuReferenceGraphExecutor::Impl +{ + bool isNodeApplicable( + const hipdnn_flatbuffers_sdk::data_objects::Node& node, + const std::unordered_map& + tensorMap) + { + try + { + auto key = buildSignatureKey(node, tensorMap); + const auto& builder = _planRegistry.getPlanBuilder(key); + return builder.isApplicable(node, tensorMap); + } + catch(...) + { + return false; + } + } + + std::unique_ptr + buildPlanForNode(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& graph, + const hipdnn_flatbuffers_sdk::data_objects::Node& node) + { + auto key = buildSignatureKey(node, graph.getTensorMap()); + + const auto& planBuilder = _planRegistry.getPlanBuilder(key); + if(!planBuilder.isApplicable(node, graph.getTensorMap())) + { + const std::string nodeName + = node.name() == nullptr ? " unknown" : " " + node.name()->str(); + throw ReferenceCapabilityError("GPU plan builder is not applicable for the given node:" + + nodeName); + } + + return planBuilder.buildNodePlan(graph, node); + } + + detail::GpuPlanBuilderRegistry _planRegistry; +}; + +GpuReferenceGraphExecutor::GpuReferenceGraphExecutor() + : _impl(std::make_unique()) +{ +} + +GpuReferenceGraphExecutor::~GpuReferenceGraphExecutor() = default; + +bool GpuReferenceGraphExecutor::isApplicable(void* graphBuffer, size_t size) +{ + auto graphWrap = hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper::fromSerializedBlob( + graphBuffer, size); + + for(uint32_t i = 0; i < graphWrap.nodeCount(); i++) + { + auto& node = graphWrap.getNode(i); + if(!_impl->isNodeApplicable(node, graphWrap.getTensorMap())) + { + return false; + } + } + return true; +} + +void GpuReferenceGraphExecutor::execute(void* graphBuffer, + size_t size, + const std::unordered_map& variantPack) +{ + auto graphWrap = hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper::fromSerializedBlob( + graphBuffer, size); + + std::vector> planExecutors; + + // The graph in graphBuffer is guaranteed to be topologically sorted. + for(uint32_t i = 0; i < graphWrap.nodeCount(); i++) + { + auto& node = graphWrap.getNode(i); + planExecutors.push_back(_impl->buildPlanForNode(graphWrap, node)); + } + + std::vector> virtualTensors; + const auto variantPackWithVirtualTensors = populateVariantPackWithMissingVirtualTensors( + variantPack, graphWrap.getTensorMap(), virtualTensors); + + for(auto& executor : planExecutors) + { + executor->execute(variantPackWithVirtualTensors); + } +} + +} // namespace hipdnn_integration_tests::gpu_graph_executor diff --git a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp index a118ebb91135..1872cb85b3cf 100644 --- a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp +++ b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp @@ -3,64 +3,37 @@ #pragma once -#include -#include +#include +#include +#include +#include -#include "detail/GpuPlanBuilderRegistry.hpp" #include "harness/IReferenceGraphExecutor.hpp" #include "harness/ReferenceCapabilityError.hpp" namespace hipdnn_integration_tests::gpu_graph_executor { +// Runs a serialized graph node by node through the GPU plan builders. The plan +// registry and the flatbuffer machinery sit behind a pimpl so they are compiled +// once (GpuReferenceGraphExecutor.cpp) instead of in every translation unit that +// needs a GPU reference. class GpuReferenceGraphExecutor : public IReferenceGraphExecutor { public: - GpuReferenceGraphExecutor() = default; + GpuReferenceGraphExecutor(); + ~GpuReferenceGraphExecutor() override; - bool isApplicable(void* graphBuffer, size_t size) override - { - auto graphWrap - = hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper::fromSerializedBlob( - graphBuffer, size); + GpuReferenceGraphExecutor(const GpuReferenceGraphExecutor&) = delete; + GpuReferenceGraphExecutor& operator=(const GpuReferenceGraphExecutor&) = delete; + GpuReferenceGraphExecutor(GpuReferenceGraphExecutor&&) = delete; + GpuReferenceGraphExecutor& operator=(GpuReferenceGraphExecutor&&) = delete; - for(uint32_t i = 0; i < graphWrap.nodeCount(); i++) - { - auto& node = graphWrap.getNode(i); - if(!isNodeApplicable(node, graphWrap.getTensorMap())) - { - return false; - } - } - return true; - } + bool isApplicable(void* graphBuffer, size_t size) override; void execute(void* graphBuffer, size_t size, - const std::unordered_map& variantPack) override - { - auto graphWrap - = hipdnn_flatbuffers_sdk::flatbuffer_utilities::GraphWrapper::fromSerializedBlob( - graphBuffer, size); - - std::vector> planExecutors; - - // The graph in graphBuffer is guaranteed to be topologically sorted. - for(uint32_t i = 0; i < graphWrap.nodeCount(); i++) - { - auto& node = graphWrap.getNode(i); - planExecutors.push_back(buildPlanForNode(graphWrap, node)); - } - - std::vector> virtualTensors; - const auto variantPackWithVirtualTensors = populateVariantPackWithMissingVirtualTensors( - variantPack, graphWrap.getTensorMap(), virtualTensors); - - for(auto& executor : planExecutors) - { - executor->execute(variantPackWithVirtualTensors); - } - } + const std::unordered_map& variantPack) override; bool requiresDeviceMemory() const override { @@ -68,151 +41,8 @@ class GpuReferenceGraphExecutor : public IReferenceGraphExecutor } private: - static std::unordered_map populateVariantPackWithMissingVirtualTensors( - const std::unordered_map& variantPack, - const std::unordered_map& - tensorMap, - std::vector>& virtualTensors) - { - auto updatedVariantPack = variantPack; - - for(const auto& [id, attr] : tensorMap) - { - if(attr->virtual_() && updatedVariantPack.find(id) == updatedVariantPack.end()) - { - auto tensor = hipdnn_test_sdk::detail::createTensorFromAttribute(*attr); - virtualTensors.push_back(std::move(tensor)); - - updatedVariantPack[id] = virtualTensors.back()->rawDeviceData(); - } - } - return updatedVariantPack; - } - - bool isNodeApplicable( - const hipdnn_flatbuffers_sdk::data_objects::Node& node, - const std::unordered_map& - tensorMap) - { - try - { - auto key = buildSignatureKey(node, tensorMap); - const auto& builder = _planRegistry.getPlanBuilder(key); - return builder.isApplicable(node, tensorMap); - } - catch(...) - { - return false; - } - } - - std::unique_ptr - buildPlanForNode(const hipdnn_flatbuffers_sdk::flatbuffer_utilities::IGraph& graph, - const hipdnn_flatbuffers_sdk::data_objects::Node& node) - { - auto key = buildSignatureKey(node, graph.getTensorMap()); - - const auto& planBuilder = _planRegistry.getPlanBuilder(key); - if(!planBuilder.isApplicable(node, graph.getTensorMap())) - { - const std::string nodeName - = node.name() == nullptr ? " unknown" : " " + node.name()->str(); - throw ReferenceCapabilityError("GPU plan builder is not applicable for the given node:" - + nodeName); - } - - return planBuilder.buildNodePlan(graph, node); - } - - static detail::GpuPlanRegistrySignatureKey buildSignatureKey( - const hipdnn_flatbuffers_sdk::data_objects::Node& node, - const std::unordered_map& - tensorMap) - { - using NodeAttrs = hipdnn_flatbuffers_sdk::data_objects::NodeAttributes; - - switch(node.attributes_type()) - { - case NodeAttrs::BatchnormAttributes: - return detail::GpuBatchnormFwdTrainSignatureKey( - node, tensorMap, node.compute_data_type()); - case NodeAttrs::ConvolutionFwdAttributes: - return detail::GpuConvolutionFwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::LayernormAttributes: - return detail::GpuLayernormFwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::LayernormBackwardAttributes: - return detail::GpuLayernormBwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::MatmulAttributes: - return detail::GpuMatmulSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::PointwiseAttributes: - return detail::GpuPointwiseSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::ReductionAttributes: - return detail::GpuReductionSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::ResampleFwdAttributes: - return detail::GpuResampleFwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::ResampleBwdAttributes: - return detail::GpuResampleBwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::RMSNormAttributes: - return detail::GpuRMSNormFwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::RMSNormBackwardAttributes: - return detail::GpuRMSNormBwdSignatureKey(node, tensorMap, node.compute_data_type()); - case NodeAttrs::BatchnormInferenceAttributes: - return detail::GpuBatchnormFwdInfSignatureKey( - node, tensorMap, node.compute_data_type()); - case NodeAttrs::BatchnormInferenceAttributesVarianceExt: - return detail::GpuBatchnormFwdInfVarianceSignatureKey( - node, tensorMap, node.compute_data_type()); - case NodeAttrs::SdpaAttributes: - { - // A ragged SDPA node (RFC-0014 packed layout) has a ragged_offset on a primary. Any - // ragged primary picks the ragged bucket, so a partly ragged node is rejected there - // instead of being run as dense. - const auto* sdpaAttributes = node.attributes_as_SdpaAttributes(); - if(sdpaAttributes != nullptr) - { - for(const auto uid : {sdpaAttributes->q_tensor_uid(), - sdpaAttributes->k_tensor_uid(), - sdpaAttributes->v_tensor_uid(), - sdpaAttributes->o_tensor_uid()}) - { - const auto it = tensorMap.find(uid); - if(it != tensorMap.end() && it->second != nullptr - && it->second->ragged_offset_tensor_uid().has_value()) - { - return detail::GpuSdpaRaggedFwdSignatureKey(node, tensorMap); - } - } - } - return detail::GpuSdpaFwdSignatureKey(node, tensorMap); - } - - // Node types with no GPU plan yet - throw descriptive error - case NodeAttrs::BatchnormBackwardAttributes: - case NodeAttrs::ConvolutionBwdAttributes: - case NodeAttrs::ConvolutionWrwAttributes: - case NodeAttrs::SdpaBackwardAttributes: - case NodeAttrs::BlockScaleDequantizeAttributes: - case NodeAttrs::BlockScaleQuantizeAttributes: - { - const std::string nodeName = node.name() == nullptr ? "unknown" : node.name()->str(); - throw ReferenceCapabilityError("GPU plan not yet implemented for node '" + nodeName - + "'. Register a GPU plan for this operation type."); - } - - case NodeAttrs::CustomOpAttributes: - throw ReferenceCapabilityError( - "GPU reference executor does not support custom operations"); - - default: - throw ReferenceCapabilityError( - "Unsupported node type for GPU signature key generation"); - } - } - - detail::GpuPlanBuilderRegistry _planRegistry; + struct Impl; + std::unique_ptr _impl; }; } // namespace hipdnn_integration_tests::gpu_graph_executor diff --git a/dnn-providers/integration-tests/tests/CMakeLists.txt b/dnn-providers/integration-tests/tests/CMakeLists.txt index 2fc735bc4cf5..a09e23bf26eb 100644 --- a/dnn-providers/integration-tests/tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/tests/CMakeLists.txt @@ -4,8 +4,9 @@ # NOTE: FrontendGraphEngineRunner.cpp, ReferenceExecutorPool.cpp and # HarnessDependencies.cpp are deliberately absent. They are the only places that # touch the frontend Graph or a hipdnnHandle_t, or hand the harness a pooled reference -# executor. CpuReferenceGraphExecutorAdapter.cpp and ReferenceGraphExecutorFactory.cpp -# ARE linked: the direct executor tests construct reference executors themselves. +# executor. CpuReferenceGraphExecutorAdapter.cpp, GpuReferenceGraphExecutor.cpp and +# ReferenceGraphExecutorFactory.cpp ARE linked: the direct executor tests construct +# reference executors themselves. # Leaving the pool, HarnessDependencies.cpp and GoldenDataRegistration.cpp out keeps # "harness code paths cannot get a pooled executor" a link-time guarantee rather than # a convention. GoldenDataRegistration.cpp is absent for the same reason: it registers @@ -25,6 +26,7 @@ add_executable(hipdnn_integration_tests_unit_tests ../src/harness/bundle/SupportVerdict.cpp ../src/harness/bundle/SweepManifestCache.cpp ../src/harness/bundle/VariantPackBuilder.cpp + ../src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp ../src/harness/input-init/FillInputs.cpp ../src/harness/input-init/SdpaBackwardInputs.cpp ../src/harness/reference-validation/BundleReferenceValidationHarness.cpp diff --git a/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp b/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp index 5502f88a8cd2..f10cdb5b66d8 100644 --- a/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp +++ b/dnn-providers/integration-tests/tests/TestGpuReferenceGraphExecutor.cpp @@ -21,6 +21,7 @@ #include #include #include +#include #include #include "ConvolutionFwdGraphTestUtils.hpp" diff --git a/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp b/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp index 05f334402b11..10d96a7fcdfe 100644 --- a/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp +++ b/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp @@ -4,6 +4,7 @@ #include #include +#include #include #include From beb6fcc5917da82f1b941617861197cc8e308917 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 09:16:51 -0600 Subject: [PATCH 05/16] refactor(hip-kernel-provider): compile the CPU reference call of the integration harness once --- .../src/integration_tests/CMakeLists.txt | 1 + .../IntegrationGraphVerificationHarness.hpp | 12 ++++++++++-- .../IntegrationGraphVerificationReference.cpp | 18 ++++++++++++++++++ 3 files changed, 29 insertions(+), 2 deletions(-) create mode 100644 dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationReference.cpp diff --git a/dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt b/dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt index cde158e4927a..c437e5e982f3 100644 --- a/dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt +++ b/dnn-providers/hip-kernel-provider/src/integration_tests/CMakeLists.txt @@ -15,6 +15,7 @@ endif() add_executable( hip_kernel_provider_integration_tests main.cpp + IntegrationGraphVerificationReference.cpp ) target_compile_options(hip_kernel_provider_integration_tests PRIVATE ${HIPDNN_WARNING_COMPILE_OPTIONS}) diff --git a/dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationHarness.hpp b/dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationHarness.hpp index 9449d35fbf47..cdeaa14b0fe4 100644 --- a/dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationHarness.hpp +++ b/dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationHarness.hpp @@ -13,9 +13,10 @@ #include #include #include -#include #include +#include +#include #include #include #include @@ -76,6 +77,13 @@ struct VerificationRun bool referenceReady = false; }; +/// Runs the CPU reference executor on a serialized graph. Defined in +/// IntegrationGraphVerificationReference.cpp so the executor's templates are compiled +/// once per executable rather than in every translation unit that includes this header. +void executeCpuReferenceGraph(void* graphBuffer, + size_t size, + const std::unordered_map& variantPack); + // NOLINTBEGIN (portability-template-virtual-member-function) template class IntegrationGraphVerificationHarness : public ::testing::TestWithParam @@ -412,7 +420,7 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam + +#include "IntegrationGraphVerificationHarness.hpp" + +namespace hip_kernel_provider::test_utilities +{ + +void executeCpuReferenceGraph(void* graphBuffer, + size_t size, + const std::unordered_map& variantPack) +{ + hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute(graphBuffer, size, variantPack); +} + +} // namespace hip_kernel_provider::test_utilities From c646bfd8691bec9fc3082f67c4829a4b20a70907 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 10:11:18 -0600 Subject: [PATCH 06/16] style(integration-tests): name the GPU executor Impl member per clang-tidy --- .../gpu-graph-executor/GpuReferenceGraphExecutor.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp index 082447ae6379..0f4bc12b3344 100644 --- a/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp +++ b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp @@ -140,7 +140,7 @@ struct GpuReferenceGraphExecutor::Impl try { auto key = buildSignatureKey(node, tensorMap); - const auto& builder = _planRegistry.getPlanBuilder(key); + const auto& builder = planRegistry.getPlanBuilder(key); return builder.isApplicable(node, tensorMap); } catch(...) @@ -155,7 +155,7 @@ struct GpuReferenceGraphExecutor::Impl { auto key = buildSignatureKey(node, graph.getTensorMap()); - const auto& planBuilder = _planRegistry.getPlanBuilder(key); + const auto& planBuilder = planRegistry.getPlanBuilder(key); if(!planBuilder.isApplicable(node, graph.getTensorMap())) { const std::string nodeName @@ -167,7 +167,7 @@ struct GpuReferenceGraphExecutor::Impl return planBuilder.buildNodePlan(graph, node); } - detail::GpuPlanBuilderRegistry _planRegistry; + detail::GpuPlanBuilderRegistry planRegistry; }; GpuReferenceGraphExecutor::GpuReferenceGraphExecutor() From b701e4bea054590a648a41bdf10c5aa39ace922b Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 12:11:34 -0600 Subject: [PATCH 07/16] ci(hipdnn): run superbuild ctest in parallel --- .github/workflows/hipdnn-superbuild-ci.yml | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index cfcaec189858..f6d73c47e27b 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -302,8 +302,9 @@ jobs: - name: Add ROCm bin to PATH run: echo "$ROCM_SDK_PATH/bin" >> "$GITHUB_PATH" + # Parallel: see the Linux lane. Serial this step was 14-20 min on Windows. - name: Run tests - run: ctest --test-dir build --output-on-failure --no-tests=error + run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 - name: Patch ROCm SDK with superbuild outputs shell: pwsh @@ -572,11 +573,14 @@ jobs: test "${#kpacks[@]}" -gt 0 || { echo "no .kpack archives produced" >&2; exit 1; } printf '%s\n' "${kpacks[@]}" + # Parallel: the suites are independent processes, and the ones that share state are + # marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. Serial this step was + # 7.7-11.3 min, dominated by a few long Python suites running back to back. - name: Run tests env: # Native rocKE AUTO detection must see the wheel COMGR used by packaging. ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} - run: ctest --test-dir build --output-on-failure --no-tests=error + run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. From 76ef8c60e1b0a2bfa262735c16ae1f9c4668890c Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 13:17:39 -0600 Subject: [PATCH 08/16] ci(hipdnn): run the full tier once, and the unlabeled tests separately --- .github/workflows/hipdnn-superbuild-ci.yml | 29 +++++++++++++++++----- 1 file changed, 23 insertions(+), 6 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index f6d73c47e27b..b2749e8e7b23 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -302,9 +302,17 @@ jobs: - name: Add ROCm bin to PATH run: echo "$ROCM_SDK_PATH/bin" >> "$GITHUB_PATH" - # Parallel: see the Linux lane. Serial this step was 14-20 min on Windows. + # Parallel: see the Linux lane for why two calls. Serial this step was 14-20 min on Windows. - name: Run tests - run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + env: + # Tests registered without the `full` label (`-L '^full$'` skips them): the ctest + # tooling tests. The golden-data suites are covered by `-L '^full$'`. + UNLABELED_TESTS: '^shared_ctest_' + run: | + rc=0 + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -R "$UNLABELED_TESTS" || rc=$? + exit $rc - name: Patch ROCm SDK with superbuild outputs shell: pwsh @@ -573,14 +581,23 @@ jobs: test "${#kpacks[@]}" -gt 0 || { echo "no .kpack archives produced" >&2; exit 1; } printf '%s\n' "${kpacks[@]}" - # Parallel: the suites are independent processes, and the ones that share state are - # marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. Serial this step was - # 7.7-11.3 min, dominated by a few long Python suites running back to back. + # `-L '^full$'` selects the unfiltered set without running the same gtest filter several + # times: the provider quick/standard/comprehensive/ffm-quick suites repeat one filter and + # `full` runs a superset, so each of them is dropped. A few tests carry no `full` label, + # so they run in the second call. Parallel: the suites are independent processes, and the + # ones that share state are marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. + # Serial this step was 7.7-11.3 min, dominated by a few long Python suites. - name: Run tests env: # Native rocKE AUTO detection must see the wheel COMGR used by packaging. ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} - run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + # Registered without the `full` label (see the comment above). + UNLABELED_TESTS: '^(shared_ctest_|hip_kernel_provider_asm_sdpa_gpu_ref_integration_tests$)' + run: | + rc=0 + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -R "$UNLABELED_TESTS" || rc=$? + exit $rc # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. From 13d5a4d4a7c5cf6f35dac2df2421f8e5faf1d789 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 14:08:07 -0600 Subject: [PATCH 09/16] Revert "ci(hipdnn): run the full tier once, and the unlabeled tests separately" This reverts commit 76ef8c60e1b0a2bfa262735c16ae1f9c4668890c. --- .github/workflows/hipdnn-superbuild-ci.yml | 29 +++++----------------- 1 file changed, 6 insertions(+), 23 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index b2749e8e7b23..f6d73c47e27b 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -302,17 +302,9 @@ jobs: - name: Add ROCm bin to PATH run: echo "$ROCM_SDK_PATH/bin" >> "$GITHUB_PATH" - # Parallel: see the Linux lane for why two calls. Serial this step was 14-20 min on Windows. + # Parallel: see the Linux lane. Serial this step was 14-20 min on Windows. - name: Run tests - env: - # Tests registered without the `full` label (`-L '^full$'` skips them): the ctest - # tooling tests. The golden-data suites are covered by `-L '^full$'`. - UNLABELED_TESTS: '^shared_ctest_' - run: | - rc=0 - ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? - ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -R "$UNLABELED_TESTS" || rc=$? - exit $rc + run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 - name: Patch ROCm SDK with superbuild outputs shell: pwsh @@ -581,23 +573,14 @@ jobs: test "${#kpacks[@]}" -gt 0 || { echo "no .kpack archives produced" >&2; exit 1; } printf '%s\n' "${kpacks[@]}" - # `-L '^full$'` selects the unfiltered set without running the same gtest filter several - # times: the provider quick/standard/comprehensive/ffm-quick suites repeat one filter and - # `full` runs a superset, so each of them is dropped. A few tests carry no `full` label, - # so they run in the second call. Parallel: the suites are independent processes, and the - # ones that share state are marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. - # Serial this step was 7.7-11.3 min, dominated by a few long Python suites. + # Parallel: the suites are independent processes, and the ones that share state are + # marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. Serial this step was + # 7.7-11.3 min, dominated by a few long Python suites running back to back. - name: Run tests env: # Native rocKE AUTO detection must see the wheel COMGR used by packaging. ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} - # Registered without the `full` label (see the comment above). - UNLABELED_TESTS: '^(shared_ctest_|hip_kernel_provider_asm_sdpa_gpu_ref_integration_tests$)' - run: | - rc=0 - ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? - ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -R "$UNLABELED_TESTS" || rc=$? - exit $rc + run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. From eeb145e9d82434ef7fa288fe5509a6e740395a3c Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 14:55:44 -0600 Subject: [PATCH 10/16] ci(hipdnn): check out shared/ctest in the superbuild lanes, report the test selection --- .github/workflows/hipdnn-superbuild-ci.yml | 29 ++++++++++++++++++++++ 1 file changed, 29 insertions(+) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index f6d73c47e27b..5a833cb4550d 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -115,6 +115,7 @@ jobs: projects/hipdnn projects/miopen test + shared/ctest # Shared CI constants, kpack-ref among them, defined in one place. - name: Load CI environment @@ -306,6 +307,19 @@ jobs: - name: Run tests run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + # TEMPORARY (this PR only, remove before merge): see the Linux lane. + - name: Report ctest selection + if: ${{ !cancelled() }} + run: | + tmp=$(mktemp -d) + ctest --test-dir build -N | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/all.txt" + ctest --test-dir build -N -L '^full$' | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/full.txt" + echo "registered: $(wc -l < "$tmp/all.txt") with label full: $(wc -l < "$tmp/full.txt")" + echo "--- without label full:" + comm -23 "$tmp/all.txt" "$tmp/full.txt" + echo "--- identical command lines (count, command):" + ctest --test-dir build -N -V | grep -E "Test command:" | sed -E 's/^[0-9]+: //' | sort | uniq -c | awk '$1>1' | cut -c1-220 + - name: Patch ROCm SDK with superbuild outputs shell: pwsh run: | @@ -375,6 +389,7 @@ jobs: projects/hipdnn projects/miopen test + shared/ctest # Shared CI constants, kpack-ref among them, defined in one place. - name: Load CI environment @@ -582,6 +597,20 @@ jobs: ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + # TEMPORARY (this PR only, remove before merge): show which tests carry the `full` label + # and which entries run an identical command line, now that shared/ctest is checked out. + - name: Report ctest selection + if: ${{ !cancelled() }} + run: | + tmp=$(mktemp -d) + ctest --test-dir build -N | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/all.txt" + ctest --test-dir build -N -L '^full$' | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/full.txt" + echo "registered: $(wc -l < "$tmp/all.txt") with label full: $(wc -l < "$tmp/full.txt")" + echo "--- without label full:" + comm -23 "$tmp/all.txt" "$tmp/full.txt" + echo "--- identical command lines (count, command):" + ctest --test-dir build -N -V | grep -E "Test command:" | sed -E 's/^[0-9]+: //' | sort | uniq -c | awk '$1>1' | cut -c1-220 + # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. # Both renderers read the committed support-claim sidecars, not the From 52dedc7761fc82d13607bc09a62f40c6acd2a64c Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 15:25:38 -0600 Subject: [PATCH 11/16] Revert "ci(hipdnn): check out shared/ctest in the superbuild lanes, report the test selection" This reverts commit eeb145e9d82434ef7fa288fe5509a6e740395a3c. --- .github/workflows/hipdnn-superbuild-ci.yml | 29 ---------------------- 1 file changed, 29 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index 5a833cb4550d..f6d73c47e27b 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -115,7 +115,6 @@ jobs: projects/hipdnn projects/miopen test - shared/ctest # Shared CI constants, kpack-ref among them, defined in one place. - name: Load CI environment @@ -307,19 +306,6 @@ jobs: - name: Run tests run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 - # TEMPORARY (this PR only, remove before merge): see the Linux lane. - - name: Report ctest selection - if: ${{ !cancelled() }} - run: | - tmp=$(mktemp -d) - ctest --test-dir build -N | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/all.txt" - ctest --test-dir build -N -L '^full$' | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/full.txt" - echo "registered: $(wc -l < "$tmp/all.txt") with label full: $(wc -l < "$tmp/full.txt")" - echo "--- without label full:" - comm -23 "$tmp/all.txt" "$tmp/full.txt" - echo "--- identical command lines (count, command):" - ctest --test-dir build -N -V | grep -E "Test command:" | sed -E 's/^[0-9]+: //' | sort | uniq -c | awk '$1>1' | cut -c1-220 - - name: Patch ROCm SDK with superbuild outputs shell: pwsh run: | @@ -389,7 +375,6 @@ jobs: projects/hipdnn projects/miopen test - shared/ctest # Shared CI constants, kpack-ref among them, defined in one place. - name: Load CI environment @@ -597,20 +582,6 @@ jobs: ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 - # TEMPORARY (this PR only, remove before merge): show which tests carry the `full` label - # and which entries run an identical command line, now that shared/ctest is checked out. - - name: Report ctest selection - if: ${{ !cancelled() }} - run: | - tmp=$(mktemp -d) - ctest --test-dir build -N | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/all.txt" - ctest --test-dir build -N -L '^full$' | sed -n 's/.*Test *#[0-9]*: //p' | sort > "$tmp/full.txt" - echo "registered: $(wc -l < "$tmp/all.txt") with label full: $(wc -l < "$tmp/full.txt")" - echo "--- without label full:" - comm -23 "$tmp/all.txt" "$tmp/full.txt" - echo "--- identical command lines (count, command):" - ctest --test-dir build -N -V | grep -E "Test command:" | sed -E 's/^[0-9]+: //' | sort | uniq -c | awk '$1>1' | cut -c1-220 - # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. # Both renderers read the committed support-claim sidecars, not the From 00a4383ae8bf5db4f62f12e885db31c76a023de8 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 15:32:27 -0600 Subject: [PATCH 12/16] ci(hipdnn): check out shared/ctest and run each gtest tier once --- .github/workflows/hipdnn-superbuild-ci.yml | 37 ++++++++++++++++++---- 1 file changed, 31 insertions(+), 6 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index f6d73c47e27b..ae626ec47833 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -114,6 +114,7 @@ jobs: dnn-providers projects/hipdnn projects/miopen + shared/ctest test # Shared CI constants, kpack-ref among them, defined in one place. @@ -302,9 +303,19 @@ jobs: - name: Add ROCm bin to PATH run: echo "$ROCM_SDK_PATH/bin" >> "$GITHUB_PATH" - # Parallel: see the Linux lane. Serial this step was 14-20 min on Windows. + # See the Linux lane for why there are two ctest calls. Serial this step was 14-20 min + # on Windows. - name: Run tests - run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + env: + # shared_ctest_optional_args_scope configures a nested CMake project whose + # find_package(Python3) must find the venv with PyYAML, not the runner's tool-cache + # Python. + VIRTUAL_ENV: ${{ env.VENV_PATH }} + run: | + rc=0 + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -LE '^(quick|standard|comprehensive|ffm-quick|full)$' || rc=$? + exit $rc - name: Patch ROCm SDK with superbuild outputs shell: pwsh @@ -374,6 +385,7 @@ jobs: dnn-providers projects/hipdnn projects/miopen + shared/ctest test # Shared CI constants, kpack-ref among them, defined in one place. @@ -573,14 +585,27 @@ jobs: test "${#kpacks[@]}" -gt 0 || { echo "no .kpack archives produced" >&2; exit 1; } printf '%s\n' "${kpacks[@]}" - # Parallel: the suites are independent processes, and the ones that share state are - # marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. Serial this step was - # 7.7-11.3 min, dominated by a few long Python suites running back to back. + # ctest runs in parallel: the suites are independent processes, and the ones that share + # state are marked RUN_SERIAL / FIXTURES_* in CMake, which ctest honors. Serial this step + # was 7.7-11.3 min, dominated by a few long Python suites running back to back. + # + # shared/ctest registers each gtest binary once per tier (quick, standard, comprehensive, + # ffm-quick, full) with the same or a narrower filter, and `full` runs a superset, so the + # `full` label selects each binary once. Entries with no tier label (rocke_*, hkp-probe-*, + # shared_ctest_*, ...) are selected by the second call. Both calls always run. - name: Run tests env: # Native rocKE AUTO detection must see the wheel COMGR used by packaging. ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} - run: ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 + # The shared_ctest_* tests configure a nested CMake project: give it Ninja (the + # runner has no make) and let its find_package(Python3) find the venv with PyYAML. + CMAKE_GENERATOR: Ninja + VIRTUAL_ENV: ${{ env.VENV_PATH }} + run: | + rc=0 + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? + ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -LE '^(quick|standard|comprehensive|ffm-quick|full)$' || rc=$? + exit $rc # The full support matrix is too large to commit (only its overview is, # in dnn-providers/integration-tests/), so publish it from every run. From 4c1ecae1286d8dad6dcdf25430605b47d7498a48 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Thu, 29 Oct 2026 16:04:15 -0600 Subject: [PATCH 13/16] ci(hipdnn): point the nested shared_ctest Python lookup at the venv --- .github/workflows/hipdnn-superbuild-ci.yml | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index ae626ec47833..9ef9f0ad2463 100644 --- a/.github/workflows/hipdnn-superbuild-ci.yml +++ b/.github/workflows/hipdnn-superbuild-ci.yml @@ -308,9 +308,9 @@ jobs: - name: Run tests env: # shared_ctest_optional_args_scope configures a nested CMake project whose - # find_package(Python3) must find the venv with PyYAML, not the runner's tool-cache - # Python. - VIRTUAL_ENV: ${{ env.VENV_PATH }} + # find_package(Python3) must find the venv with PyYAML. actions/setup-python exports + # Python3_ROOT_DIR for its own Python, which has no PyYAML, so override it. + Python3_ROOT_DIR: ${{ env.VENV_PATH }} run: | rc=0 ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? @@ -598,9 +598,11 @@ jobs: # Native rocKE AUTO detection must see the wheel COMGR used by packaging. ROCKE_COMGR_LIB: ${{ env.HKP_ROCKE_COMGR_LIB }} # The shared_ctest_* tests configure a nested CMake project: give it Ninja (the - # runner has no make) and let its find_package(Python3) find the venv with PyYAML. + # runner has no make), and point its find_package(Python3) at the venv with PyYAML. + # actions/setup-python exports Python3_ROOT_DIR for its own Python, which has no + # PyYAML and takes precedence over VIRTUAL_ENV. CMAKE_GENERATOR: Ninja - VIRTUAL_ENV: ${{ env.VENV_PATH }} + Python3_ROOT_DIR: ${{ env.VENV_PATH }} run: | rc=0 ctest --test-dir build --output-on-failure --no-tests=error --parallel 16 -L '^full$' || rc=$? From 7503624ee8d3e4d94bd3ed7badaa43abeb1761f9 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Fri, 30 Oct 2026 15:29:32 -0600 Subject: [PATCH 14/16] fix(integration-tests): decide YAML test categories before registering hipdnn_golden_data_tests The golden-data binary was registered before DNN_PROVIDER_TEST_CATEGORY_YAMLS was set, so _add_test_target_internal() still added the raw, unlabeled hipdnn_golden_data_tests test next to its tiered *_suite entries and the binary ran twice. Move the categorization block above the registration. --- .../integration-tests/CMakeLists.txt | 85 ++++++++++--------- 1 file changed, 45 insertions(+), 40 deletions(-) diff --git a/dnn-providers/integration-tests/CMakeLists.txt b/dnn-providers/integration-tests/CMakeLists.txt index 047af95bb4bd..12a66a5c1713 100644 --- a/dnn-providers/integration-tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/CMakeLists.txt @@ -300,6 +300,51 @@ target_compile_options(${GOLDEN_DATA_TESTS_EXE} clang_tidy_check(${GOLDEN_DATA_TESTS_EXE}) +# YAML test categorization has to be decided before the first add_*_test_target() call below: +# _add_test_target_internal() (dnn-providers/cmake/Tests.cmake) skips the raw, unlabeled +# add_test() for a target only while DNN_PROVIDER_TEST_CATEGORY_YAMLS is already set. Set +# after hipdnn_golden_data_tests was registered, the binary ran twice: once as the tiered +# *_suite entries and once as an unlabeled raw test. +set(HIPDNN_IT_GTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/test_categories.yaml") +set(HIPDNN_IT_CTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/test_categories_external.yaml") +set(HIPDNN_IT_YAML_CATEGORIZATION_ENABLED FALSE) +if(EXISTS "${HIPDNN_IT_GTEST_CATEGORIES_YAML}" + AND EXISTS "${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake") + include("${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake") + set(DNN_PROVIDER_TEST_CATEGORY_YAMLS "${HIPDNN_IT_GTEST_CATEGORIES_YAML}") + set(DNN_PROVIDER_CTEST_CATEGORIES_YAML "${HIPDNN_IT_CTEST_CATEGORIES_YAML}") + set(HIPDNN_IT_YAML_CATEGORIZATION_ENABLED TRUE) + message(STATUS "hipdnn-integration-tests: YAML-based CTest categorization enabled") + # Label the pre-registered bundle-verifier ctest test via the shared YAML path + # (it is registered directly above, not by _create_test_name_validation_target_internal). + if(HIPDNN_IT_BUNDLE_VERIFIER_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_BUNDLE_VERIFIER_TEST}) + endif() + if(HIPDNN_IT_SUPPORT_CLAIM_VERIFIER_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_CLAIM_VERIFIER_TEST}) + endif() + if(HIPDNN_IT_SUPPORT_CLAIM_HARVEST_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_CLAIM_HARVEST_TEST}) + endif() + if(HIPDNN_IT_SUPPORT_SIDECAR_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_SIDECAR_TEST}) + endif() + if(HIPDNN_IT_BUNDLE_DISCOVERY_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_BUNDLE_DISCOVERY_TEST}) + endif() + if(HIPDNN_IT_SUPPORT_MATRIX_RENDERER_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_MATRIX_RENDERER_TEST}) + endif() + if(HIPDNN_IT_SUPPORT_FEATURES_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_FEATURES_TEST}) + endif() + if(HIPDNN_IT_GOLDEN_GENERATOR_TEST) + _apply_provider_ctest_category_labels(${HIPDNN_IT_GOLDEN_GENERATOR_TEST}) + endif() +else() + message(STATUS "hipdnn-integration-tests: shared/ctest or test_categories.yaml not found, using legacy CTest labels") +endif() + # SKIP_NAME_VALIDATION: every test this binary registers is named after a bundle's # directory path (deriveTestName() in BundleDiscovery.hpp), so the names are # lower_snake_case by construction -- `quick_SdpaFwd_bhsd_bf16_hd128_causal_batch_ @@ -341,46 +386,6 @@ clang_tidy_check(${INTEGRATION_TESTS_EXE}) # (see dnn-providers/miopen-provider) which runs this binary against a specific # engine and selects tests via that provider's test_categories_integration.yaml. -set(HIPDNN_IT_GTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/test_categories.yaml") -set(HIPDNN_IT_CTEST_CATEGORIES_YAML "${CMAKE_CURRENT_SOURCE_DIR}/test_categories_external.yaml") -set(HIPDNN_IT_YAML_CATEGORIZATION_ENABLED FALSE) -if(EXISTS "${HIPDNN_IT_GTEST_CATEGORIES_YAML}" - AND EXISTS "${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake") - include("${ROCM_LIBRARIES_ROOT}/shared/ctest/TestCategories.cmake") - set(DNN_PROVIDER_TEST_CATEGORY_YAMLS "${HIPDNN_IT_GTEST_CATEGORIES_YAML}") - set(DNN_PROVIDER_CTEST_CATEGORIES_YAML "${HIPDNN_IT_CTEST_CATEGORIES_YAML}") - set(HIPDNN_IT_YAML_CATEGORIZATION_ENABLED TRUE) - message(STATUS "hipdnn-integration-tests: YAML-based CTest categorization enabled") - # Label the pre-registered bundle-verifier ctest test via the shared YAML path - # (it is registered directly above, not by _create_test_name_validation_target_internal). - if(HIPDNN_IT_BUNDLE_VERIFIER_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_BUNDLE_VERIFIER_TEST}) - endif() - if(HIPDNN_IT_SUPPORT_CLAIM_VERIFIER_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_CLAIM_VERIFIER_TEST}) - endif() - if(HIPDNN_IT_SUPPORT_CLAIM_HARVEST_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_CLAIM_HARVEST_TEST}) - endif() - if(HIPDNN_IT_SUPPORT_SIDECAR_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_SIDECAR_TEST}) - endif() - if(HIPDNN_IT_BUNDLE_DISCOVERY_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_BUNDLE_DISCOVERY_TEST}) - endif() - if(HIPDNN_IT_SUPPORT_MATRIX_RENDERER_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_MATRIX_RENDERER_TEST}) - endif() - if(HIPDNN_IT_SUPPORT_FEATURES_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_SUPPORT_FEATURES_TEST}) - endif() - if(HIPDNN_IT_GOLDEN_GENERATOR_TEST) - _apply_provider_ctest_category_labels(${HIPDNN_IT_GOLDEN_GENERATOR_TEST}) - endif() -else() - message(STATUS "hipdnn-integration-tests: shared/ctest or test_categories.yaml not found, using legacy CTest labels") -endif() - add_subdirectory(gpu-ref) add_subdirectory(tests) From 0528e30579f82046a069473bd0dff3b7dd655c67 Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Fri, 30 Oct 2026 16:06:32 -0600 Subject: [PATCH 15/16] fix(hip-kernel-provider): place the FlyDSL test TOMLs beside the census shard hkp_reserve_census_shard() was called for each FlyDSL engine and architecture without the generated TOML, so the shard check added in #14573 failed configure for any architecture with FlyDSL content (for example gfx1151: "the gfx1151 census shard names hipkernel_flydsl_rmsnorm_gfx1151.toml, but nothing places it beside the shard"). Pass the file, as the gfx950 registration already does. --- dnn-providers/hip-kernel-provider/src/CMakeLists.txt | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/dnn-providers/hip-kernel-provider/src/CMakeLists.txt b/dnn-providers/hip-kernel-provider/src/CMakeLists.txt index 0f002ac6d86d..613525256133 100644 --- a/dnn-providers/hip-kernel-provider/src/CMakeLists.txt +++ b/dnn-providers/hip-kernel-provider/src/CMakeLists.txt @@ -491,7 +491,10 @@ if(HIPKERNELPROVIDER_ENABLE_TESTS) set_property(GLOBAL PROPERTY HKP_CENSUS_EXTERNAL_TEST_FILE_${_arch} "${_shard_ctest}") endif() - hkp_reserve_census_shard(_shard_dest ${_arch}) + # The TOML goes to the reservation as well as TEST_CONFIG, or the shard check + # in _hkp_write_census_shard() fails configure: nothing else places it + # beside the build-tree shard. + hkp_reserve_census_shard(_shard_dest ${_arch} FILES "${_toml}") set(_category_args "") if(HIPKERNELPROVIDER_YAML_CATEGORIZATION_ENABLED) From 5c41e0713bf817b9b64e126a27d3a6f85666667e Mon Sep 17 00:00:00 2001 From: Adam Dickin Date: Mon, 2 Nov 2026 16:31:00 -0600 Subject: [PATCH 16/16] fix(integration-tests): gate costly SdpaBackward goldens off the CPU reference The scalar CPU SdpaBackward reference needs seconds to tens of seconds per 256x128 bundle, and twelve of them in the quick tier exceeded the 600 s ctest timeout of hipdnn_golden_data_tests_quick_suite. referenceShapeIsAffordable now caps the CPU backward at 4M multiply-accumulates per Q*K^T pass (Q elements times Skv), in addition to the existing quick-tier limit. The 64-row padded and GQA bundles stay on CPU; the excluded bundles stay covered by the GPU lane. --- .../ReferenceOpCoverage.cpp | 61 ++++++++++++++----- .../ReferenceOpCoverage.hpp | 5 +- .../tests/ReferenceGraphFixtures.hpp | 43 +++++++++++++ .../tests/TestReferenceOpCoverage.cpp | 33 ++++++++++ 4 files changed, 126 insertions(+), 16 deletions(-) diff --git a/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.cpp b/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.cpp index a8a2e64b1883..e17821a3128d 100644 --- a/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.cpp +++ b/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.cpp @@ -116,6 +116,14 @@ bool referenceCoversGraph(ReferenceExecutorType type, const void* graphBuffer, s constexpr std::string_view K_CPU_TIER_PREFIX = "quick_"; constexpr int64_t K_CPU_MAX_WORKING_SET_ELEMENTS = int64_t{4} * 1024 * 1024; +// Backward is capped on its own axis because its cost does not follow the forward cap: the +// scalar backward needs seconds to tens of seconds at a 256x128 head (8.4M MACs per Q*K^T +// pass) while the 64-row shapes below 3.2M finish in about a second or less, so the 4M +// element cap above would admit every one of them. The threshold sits between those two +// clusters. Measured on a loaded CI runner, so only the gap, not the absolute times, is +// meaningful. +constexpr int64_t K_CPU_MAX_SDPA_BWD_MACS = int64_t{4} * 1024 * 1024; + namespace { @@ -134,10 +142,9 @@ int64_t elementCount(const hipdnn_flatbuffers_sdk::data_objects::TensorAttribute return elements; } -// Both caps apply to Sdpa only, so a graph with no Sdpa node is affordable by -// construction. Batchnorm and the other CPU-covered ops are cheap at every -// checked-in shape, and gating them would trade real coverage for no measurable -// time. +// Both caps apply to Sdpa forward and backward only, so a graph with no Sdpa node is +// affordable by construction. Batchnorm and the other CPU-covered ops are cheap at every +// checked-in shape, and gating them would trade real coverage for no measurable time. bool isSdpaShapeAffordableOnCpuRunner(std::string_view bundleId, const void* graphBuffer, size_t size) @@ -148,10 +155,18 @@ bool isSdpaShapeAffordableOnCpuRunner(std::string_view bundleId, graphBuffer, size); const auto& tensorMap = graph.getTensorMap(); + const auto lookup = [&tensorMap](int64_t uid) { + const auto it = tensorMap.find(uid); + return it == tensorMap.end() ? nullptr : it->second; + }; + for(uint32_t i = 0; i < graph.nodeCount(); ++i) { const auto& node = graph.getNode(i); - if(node.attributes_type() != NodeAttributes::SdpaAttributes) + const bool isForward = node.attributes_type() == NodeAttributes::SdpaAttributes; + const bool isBackward + = node.attributes_type() == NodeAttributes::SdpaBackwardAttributes; + if(!isForward && !isBackward) { continue; } @@ -161,20 +176,38 @@ bool isSdpaShapeAffordableOnCpuRunner(std::string_view bundleId, return false; } - const auto* attributes = node.attributes_as_SdpaAttributes(); - if(attributes == nullptr) + if(isForward) { + const auto* attributes = node.attributes_as_SdpaAttributes(); + if(attributes == nullptr) + { + continue; + } + + const int64_t working = std::max(elementCount(lookup(attributes->q_tensor_uid())), + elementCount(lookup(attributes->k_tensor_uid()))); + if(working > K_CPU_MAX_WORKING_SET_ELEMENTS) + { + return false; + } continue; } - const auto lookup = [&tensorMap](int64_t uid) { - const auto it = tensorMap.find(uid); - return it == tensorMap.end() ? nullptr : it->second; - }; + const auto* attributes = node.attributes_as_SdpaBackwardAttributes(); + if(attributes == nullptr) + { + continue; + } - const int64_t working = std::max(elementCount(lookup(attributes->q_tensor_uid())), - elementCount(lookup(attributes->k_tensor_uid()))); - if(working > K_CPU_MAX_WORKING_SET_ELEMENTS) + // Q is [B, H, Sq, D] and K is [B, Hkv, Skv, D], so Q's element count times Skv is + // the multiply-accumulate count of one Q*K^T pass (B * H * Sq * Skv * D). + const auto* k = lookup(attributes->k_tensor_uid()); + if(k == nullptr || k->dims() == nullptr || k->dims()->size() < 2) + { + continue; + } + const int64_t seqKv = (*k->dims())[k->dims()->size() - 2]; + if(elementCount(lookup(attributes->q_tensor_uid())) * seqKv > K_CPU_MAX_SDPA_BWD_MACS) { return false; } diff --git a/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.hpp b/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.hpp index 3d09f4dff3cb..48f02e01386c 100644 --- a/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.hpp +++ b/dnn-providers/integration-tests/src/harness/reference-validation/ReferenceOpCoverage.hpp @@ -59,8 +59,9 @@ bool referenceCoversGraph(ReferenceExecutorType type, const void* graphBuffer, s /// shape worth spending the time on". Folding the second into the first would make /// a runtime-cost decision look like a missing capability. /// -/// Only the CPU reference is gated, and only for Sdpa: it is scalar, so it is -/// restricted to the `quick` tier and to shapes under a working-set cap. Excluded +/// Only the CPU reference is gated, and only for Sdpa forward and backward: it is scalar, so +/// it is restricted to the `quick` tier and, per direction, to shapes under a cost cap (a +/// working-set cap for the forward, a multiply-accumulate cap for the backward). Excluded /// bundles stay covered by the GPU lane, so this trades CPU cross-checking of the /// larger shapes for a golden-data run that finishes. That trade only holds while /// the GPU lane actually runs, which is the planner's job to establish -- see diff --git a/dnn-providers/integration-tests/tests/ReferenceGraphFixtures.hpp b/dnn-providers/integration-tests/tests/ReferenceGraphFixtures.hpp index b6ff2a87c3a3..24a211e54af0 100644 --- a/dnn-providers/integration-tests/tests/ReferenceGraphFixtures.hpp +++ b/dnn-providers/integration-tests/tests/ReferenceGraphFixtures.hpp @@ -8,8 +8,10 @@ // without depending on the checked-in bundles. #include +#include #include +#include #include #include @@ -141,4 +143,45 @@ inline flatbuffers::DetachedBuffer buildSdpaGraph(int64_t seq) return builder.Release(); } +// A single-node SdpaBackward graph. Only Q and K carry the dimensions the cost gate reads +// (Q is [batch, heads, seqQ, headDim], K is [batch, heads, seqKv, headDim]); the other +// operands reuse Q's shape because the gate never looks at them. +inline flatbuffers::DetachedBuffer buildSdpaBackwardGraph( + int64_t batch, int64_t heads, int64_t seqQ, int64_t seqKv, int64_t headDim) +{ + using namespace hipdnn_flatbuffers_sdk::data_objects; + + flatbuffers::FlatBufferBuilder builder; + const auto strides = [headDim](int64_t seq, int64_t headCount) { + return std::vector{headCount * seq * headDim, seq * headDim, headDim, 1}; + }; + const std::vector qDims = {batch, heads, seqQ, headDim}; + const std::vector kDims = {batch, heads, seqKv, headDim}; + const auto qStrides = strides(seqQ, heads); + const auto kStrides = strides(seqKv, heads); + + std::vector> tensors; + tensors.push_back( + CreateTensorAttributesDirect(builder, 1, "q", DataType::HALF, &qStrides, &qDims)); + tensors.push_back( + CreateTensorAttributesDirect(builder, 2, "k", DataType::HALF, &kStrides, &kDims)); + for(int64_t uid = 3; uid <= 9; ++uid) + { + tensors.push_back( + CreateTensorAttributesDirect(builder, uid, "", DataType::HALF, &qStrides, &qDims)); + } + + auto attributes = CreateSdpaBackwardAttributes(builder, 1, 2, 3, 4, 5, 6, 7, 8, 9); + std::vector> nodes; + nodes.push_back(CreateNodeDirect(builder, + "sdpa_bwd", + DataType::FLOAT, + NodeAttributes::SdpaBackwardAttributes, + attributes.Union())); + + builder.Finish(CreateGraphDirect( + builder, "test", DataType::HALF, DataType::FLOAT, DataType::FLOAT, &tensors, &nodes)); + return builder.Release(); +} + } // namespace hipdnn_integration_tests::bundle::fixtures diff --git a/dnn-providers/integration-tests/tests/TestReferenceOpCoverage.cpp b/dnn-providers/integration-tests/tests/TestReferenceOpCoverage.cpp index b4053ce5fa95..d29d1de2cb11 100644 --- a/dnn-providers/integration-tests/tests/TestReferenceOpCoverage.cpp +++ b/dnn-providers/integration-tests/tests/TestReferenceOpCoverage.cpp @@ -35,6 +35,7 @@ using hipdnn_integration_tests::bundle::referenceShapeIsAffordable; using hipdnn_integration_tests::bundle::referenceSupportedOps; using hipdnn_integration_tests::bundle::uncoveredNodeTypes; using hipdnn_integration_tests::bundle::fixtures::buildBatchnormGraph; +using hipdnn_integration_tests::bundle::fixtures::buildSdpaBackwardGraph; using hipdnn_integration_tests::bundle::fixtures::buildSdpaGraph; // NOLINTBEGIN(readability-identifier-naming) @@ -279,6 +280,38 @@ TEST(TestReferenceOpCoverage, CpuSdpaIsLimitedByWorkingSet) ReferenceExecutorType::CPU, "quick_SdpaFwd_x.Small", hugeGraph.data(), hugeGraph.size())); } +// The scalar CPU backward costs far more than the forward at the same element count, so it +// has its own cap in multiply-accumulates. The two shapes are the checked-in quick +// SdpaBackward bundles on either side of it: the 64-row padded/GQA bundles (about a +// second) and the 256x128 head bundles (tens of seconds each, enough together to time out +// the whole quick suite). +TEST(TestReferenceOpCoverage, CpuSdpaBackwardIsLimitedByCost) +{ + const auto cheap = buildSdpaBackwardGraph(2, 4, 64, 96, 64); + const auto costly = buildSdpaBackwardGraph(1, 1, 256, 256, 128); + + EXPECT_TRUE(referenceShapeIsAffordable( + ReferenceExecutorType::CPU, "quick_SdpaBackward_x.Small", cheap.data(), cheap.size())); + EXPECT_FALSE(referenceShapeIsAffordable( + ReferenceExecutorType::CPU, "quick_SdpaBackward_x.Small", costly.data(), costly.size())); +} + +TEST(TestReferenceOpCoverage, CpuSdpaBackwardIsLimitedToQuickTier) +{ + const auto cheap = buildSdpaBackwardGraph(2, 4, 64, 96, 64); + + EXPECT_FALSE(referenceShapeIsAffordable( + ReferenceExecutorType::CPU, "standard_SdpaBackward_x.Small", cheap.data(), cheap.size())); +} + +TEST(TestReferenceOpCoverage, SdpaBackwardCostGateIsCpuOnly) +{ + const auto costly = buildSdpaBackwardGraph(1, 1, 256, 256, 128); + + EXPECT_TRUE(referenceShapeIsAffordable( + ReferenceExecutorType::GPU, "quick_SdpaBackward_x.Small", costly.data(), costly.size())); +} + // The gate is CPU-only: the GPU reference runs every one of these in milliseconds // and is what keeps the excluded bundles covered. TEST(TestReferenceOpCoverage, OnlyTheCpuReferenceIsGatedOnCost)