diff --git a/.github/workflows/hipdnn-superbuild-ci.yml b/.github/workflows/hipdnn-superbuild-ci.yml index 5a8cf4fef9d2..1e0a5199ead5 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. @@ -303,8 +304,19 @@ jobs: - name: Add ROCm bin to PATH run: echo "$ROCM_SDK_PATH/bin" >> "$GITHUB_PATH" + # 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 + env: + # shared_ctest_optional_args_scope configures a nested CMake project whose + # 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=$? + 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 +386,7 @@ jobs: dnn-providers projects/hipdnn projects/miopen + shared/ctest test # Shared CI constants, kpack-ref among them, defined in one place. @@ -574,11 +587,29 @@ jobs: test "${#kpacks[@]}" -gt 0 || { echo "no .kpack archives produced" >&2; exit 1; } printf '%s\n' "${kpacks[@]}" + # 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 + # The shared_ctest_* tests configure a nested CMake project: give it Ninja (the + # 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 + Python3_ROOT_DIR: ${{ 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. 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 diff --git a/dnn-providers/integration-tests/CMakeLists.txt b/dnn-providers/integration-tests/CMakeLists.txt index 8c787143d77d..e46a778697f8 100644 --- a/dnn-providers/integration-tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/CMakeLists.txt @@ -246,7 +246,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 @@ -259,6 +261,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 ) @@ -278,6 +281,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 @@ -286,6 +290,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 @@ -318,47 +323,11 @@ target_compile_options(${GOLDEN_DATA_TESTS_EXE} clang_tidy_check(${GOLDEN_DATA_TESTS_EXE}) -# 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_ -# Small_CpuRef.Small`. The PascalCase convention the validator enforces is for -# hand-written C++ test names; satisfying it here would mean renaming the checked-in -# bundle tree, which is the data, not the test code. hipdnn_integration_tests -# registers the same bundle-derived names and is likewise absent from the validated -# set, because it is not a tiered ctest target (see the note below). -add_integration_test_target(${GOLDEN_DATA_TESTS_EXE} ${CMAKE_CURRENT_BINARY_DIR} - SKIP_NAME_VALIDATION) - -add_subdirectory(src/integration-tests) - -target_include_directories(${INTEGRATION_TESTS_EXE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/src) -target_include_directories(${INTEGRATION_TESTS_EXE} SYSTEM PRIVATE - ${CMAKE_CURRENT_SOURCE_DIR}/src/third-party/argparse - ${CMAKE_CURRENT_SOURCE_DIR}/src/third-party/tomlplusplus -) - -target_link_libraries(${INTEGRATION_TESTS_EXE} - PRIVATE - hip::host - hipdnn_gpu_ref - hipdnn_test_sdk - hipdnn_plugin_sdk - GTest::gtest - Threads::Threads - $<$:roc::rocrand> - $<$:shlwapi> -) - -target_compile_options(${INTEGRATION_TESTS_EXE} PRIVATE ${INTEGRATION_TESTS_WARNING_COMPILE_OPTIONS}) - -clang_tidy_check(${INTEGRATION_TESTS_EXE}) - -# NOTE: hipdnn_integration_tests is intentionally NOT registered as a tiered -# ctest target here. The bundle sweep tests it carries exercise the plugin, so -# they are driven by each provider's add_external_integration_test_target() -# (see dnn-providers/miopen-provider) which runs this binary against a specific -# engine and selects tests via that provider's test_categories_integration.yaml. - +# 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) @@ -405,6 +374,47 @@ 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_ +# Small_CpuRef.Small`. The PascalCase convention the validator enforces is for +# hand-written C++ test names; satisfying it here would mean renaming the checked-in +# bundle tree, which is the data, not the test code. hipdnn_integration_tests +# registers the same bundle-derived names and is likewise absent from the validated +# set, because it is not a tiered ctest target (see the note below). +add_integration_test_target(${GOLDEN_DATA_TESTS_EXE} ${CMAKE_CURRENT_BINARY_DIR} + SKIP_NAME_VALIDATION) + +add_subdirectory(src/integration-tests) + +target_include_directories(${INTEGRATION_TESTS_EXE} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/src) +target_include_directories(${INTEGRATION_TESTS_EXE} SYSTEM PRIVATE + ${CMAKE_CURRENT_SOURCE_DIR}/src/third-party/argparse + ${CMAKE_CURRENT_SOURCE_DIR}/src/third-party/tomlplusplus +) + +target_link_libraries(${INTEGRATION_TESTS_EXE} + PRIVATE + hip::host + hipdnn_gpu_ref + hipdnn_test_sdk + hipdnn_plugin_sdk + GTest::gtest + Threads::Threads + $<$:roc::rocrand> + $<$:shlwapi> +) + +target_compile_options(${INTEGRATION_TESTS_EXE} PRIVATE ${INTEGRATION_TESTS_WARNING_COMPILE_OPTIONS}) + +clang_tidy_check(${INTEGRATION_TESTS_EXE}) + +# NOTE: hipdnn_integration_tests is intentionally NOT registered as a tiered +# ctest target here. The bundle sweep tests it carries exercise the plugin, so +# they are driven by each provider's add_external_integration_test_target() +# (see dnn-providers/miopen-provider) which runs this binary against a specific +# engine and selects tests via that provider's test_categories_integration.yaml. + add_subdirectory(gpu-ref) add_subdirectory(tests) 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/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/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/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp new file mode 100644 index 000000000000..397f40546363 --- /dev/null +++ b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/GpuReferenceGraphExecutor.cpp @@ -0,0 +1,400 @@ +// 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 +#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; +} + +// Owns one device allocation; frees it after the device is idle. +struct DeviceBuffer +{ + void* ptr = nullptr; + // Graph tensors have positive dimensions and empty sequence-length lists are not + // bound, so a zero-byte buffer is a caller bug. + explicit DeviceBuffer(size_t bytes) + { + if(bytes == 0) + { + throw std::invalid_argument("ragged SDPA densify: empty device buffer"); + } + hipdnn_gpu_ref::detail::throwOnHipError(hipMalloc(&ptr, bytes), + "ragged SDPA densify: hipMalloc"); + } + DeviceBuffer(const DeviceBuffer&) = delete; + DeviceBuffer& operator=(const DeviceBuffer&) = delete; + ~DeviceBuffer() + { + (void)hipDeviceSynchronize(); + (void)hipFree(ptr); + } +}; + +void copyBytes(void* to, const void* from, size_t bytes, hipMemcpyKind kind) +{ + if(bytes > 0) + { + hipdnn_gpu_ref::detail::throwOnHipError(hipMemcpy(to, from, bytes, kind), + "ragged SDPA densify: hipMemcpy"); + } +} + +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 node with any ragged primary is keyed as ragged. A partly ragged node is declined + // by both plans (the ragged plan needs all four primaries ragged, and the dense plan + // rejects any ragged tensor); isApplicable then tries the dense rewrite. + const auto* sdpaAttributes = node.attributes_as_SdpaAttributes(); + if(sdpaAttributes != nullptr + && hipdnn_test_sdk::detail::hasRaggedSdpaPrimary(*sdpaAttributes, tensorMap)) + { + return detail::GpuSdpaRaggedFwdSignatureKey(node, tensorMap); + } + return detail::GpuSdpaFwdSignatureKey(node, tensorMap); + } + case NodeAttrs::SdpaBackwardAttributes: + return detail::GpuSdpaBwdSignatureKey(node, tensorMap); + + // Node types with no GPU plan yet - throw descriptive error + case NodeAttrs::BatchnormBackwardAttributes: + case NodeAttrs::ConvolutionBwdAttributes: + case NodeAttrs::ConvolutionWrwAttributes: + 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; + } + } + + bool isApplicableAsIs(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(!isNodeApplicable(node, graphWrap.getTensorMap())) + { + return false; + } + } + return true; + } + + // A ragged SDPA graph no plan accepts as is runs as its dense padded equivalent + // (RaggedSdpaDensify.hpp), as in the CPU executor. + bool isApplicable(void* graphBuffer, size_t size) + { + if(isApplicableAsIs(graphBuffer, size)) + { + return true; + } + auto dense = hipdnn_test_sdk::detail::densifyRaggedSdpaGraph(graphBuffer, size); + return dense.has_value() && isApplicableAsIs(dense->graph.data(), dense->graph.size()); + } + + void execute(void* graphBuffer, + size_t size, + const std::unordered_map& variantPack) + { + if(!isApplicableAsIs(graphBuffer, size)) + { + if(auto dense = hipdnn_test_sdk::detail::densifyRaggedSdpaGraph(graphBuffer, size); + dense.has_value()) + { + executeRaggedAsDense(*dense, variantPack); + return; + } + } + executeAsIs(graphBuffer, size, variantPack); + } + + void executeAsIs(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(buildPlanForNode(graphWrap, node)); + } + + std::vector> virtualTensors; + const auto variantPackWithVirtualTensors = populateVariantPackWithMissingVirtualTensors( + variantPack, graphWrap.getTensorMap(), virtualTensors); + + for(auto& executor : planExecutors) + { + executor->execute(variantPackWithVirtualTensors); + } + } + + // The CPU executor's ragged-as-dense run with device buffers: offsets and packed inputs are + // read back, unpacked on the host and uploaded; dense outputs are read back, packed into a + // host copy of the packed output (so bytes outside live tokens are kept) and uploaded. + void executeRaggedAsDense(hipdnn_test_sdk::detail::DenseSdpaFromRagged& dense, + const std::unordered_map& variantPack) + { + namespace rd = hipdnn_test_sdk::detail; + const auto& primaries = dense.primaries; + std::vector> offsets; + std::vector> lengths; + for(const auto& primary : primaries) + { + std::vector table(static_cast(dense.batch + 1) + * primary.offsetElementSize); + copyBytes(table.data(), + variantPack.at(primary.offsetUid), + table.size(), + hipMemcpyDeviceToHost); + offsets.push_back(rd::raggedElementOffsets(primary, table.data(), dense.batch)); + lengths.push_back(rd::raggedLengths(primary, offsets.back())); + } + const auto [qLengths, kvLengths] = rd::raggedGroupLengths(dense, lengths); + + auto densePack = variantPack; + std::vector> deviceBuffers; + std::vector> hostDense(primaries.size()); + for(size_t i = 0; i < primaries.size(); ++i) + { + hostDense[i].assign(rd::raggedDenseBytes(primaries[i]), 0); + if(!primaries[i].output) + { + std::vector packed( + rd::raggedPackedBytes(primaries[i], offsets[i], lengths[i])); + copyBytes(packed.data(), + variantPack.at(primaries[i].uid), + packed.size(), + hipMemcpyDeviceToHost); + rd::copyRaggedTokens( + primaries[i], offsets[i], lengths[i], packed.data(), hostDense[i].data(), true); + } + deviceBuffers.push_back(std::make_unique(hostDense[i].size())); + copyBytes(deviceBuffers.back()->ptr, + hostDense[i].data(), + hostDense[i].size(), + hipMemcpyHostToDevice); + densePack[primaries[i].uid] = deviceBuffers.back()->ptr; + } + const auto bindSeqLens = [&](std::optional added, + int64_t uid, + const std::vector& packed, + const char* name) { + if(packed.empty()) + { + return; + } + std::vector values(packed.size()); + if(added.has_value()) + { + values.assign(packed.begin(), packed.end()); + deviceBuffers.push_back( + std::make_unique(values.size() * sizeof(int32_t))); + copyBytes(deviceBuffers.back()->ptr, + values.data(), + values.size() * sizeof(int32_t), + hipMemcpyHostToDevice); + densePack[uid] = deviceBuffers.back()->ptr; + } + else + { + copyBytes(values.data(), + variantPack.at(uid), + values.size() * sizeof(int32_t), + hipMemcpyDeviceToHost); + rd::requireSeqLensWithinPacked(values.data(), packed, name); + } + }; + bindSeqLens(dense.addedSeqLenQUid, dense.seqLenQUid, qLengths, "seq_len_q"); + bindSeqLens(dense.addedSeqLenKvUid, dense.seqLenKvUid, kvLengths, "seq_len_kv"); + + executeAsIs(dense.graph.data(), dense.graph.size(), densePack); + + for(size_t i = 0; i < primaries.size(); ++i) + { + if(!primaries[i].output) + { + continue; + } + copyBytes(hostDense[i].data(), + densePack.at(primaries[i].uid), + hostDense[i].size(), + hipMemcpyDeviceToHost); + std::vector packed( + rd::raggedPackedBytes(primaries[i], offsets[i], lengths[i])); + copyBytes(packed.data(), + variantPack.at(primaries[i].uid), + packed.size(), + hipMemcpyDeviceToHost); + rd::copyRaggedTokens( + primaries[i], offsets[i], lengths[i], packed.data(), hostDense[i].data(), false); + copyBytes(variantPack.at(primaries[i].uid), + packed.data(), + packed.size(), + hipMemcpyHostToDevice); + } + } + + 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) +{ + return _impl->isApplicable(graphBuffer, size); +} + +void GpuReferenceGraphExecutor::execute(void* graphBuffer, + size_t size, + const std::unordered_map& variantPack) +{ + _impl->execute(graphBuffer, size, variantPack); +} + +} // 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 a8df88bac7c8..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,58 +3,37 @@ #pragma once +#include #include #include -#include -#include -#include +#include -#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; - // A ragged SDPA graph no plan accepts as is runs as its dense padded equivalent - // (RaggedSdpaDensify.hpp), as in the CPU executor. - bool isApplicable(void* graphBuffer, size_t size) override - { - if(isApplicableAsIs(graphBuffer, size)) - { - return true; - } - auto dense = hipdnn_test_sdk::detail::densifyRaggedSdpaGraph(graphBuffer, size); - return dense.has_value() && isApplicableAsIs(dense->graph.data(), dense->graph.size()); - } + GpuReferenceGraphExecutor(const GpuReferenceGraphExecutor&) = delete; + GpuReferenceGraphExecutor& operator=(const GpuReferenceGraphExecutor&) = delete; + GpuReferenceGraphExecutor(GpuReferenceGraphExecutor&&) = delete; + GpuReferenceGraphExecutor& operator=(GpuReferenceGraphExecutor&&) = delete; + + bool isApplicable(void* graphBuffer, size_t size) override; void execute(void* graphBuffer, size_t size, - const std::unordered_map& variantPack) override - { - if(!isApplicableAsIs(graphBuffer, size)) - { - if(auto dense = hipdnn_test_sdk::detail::densifyRaggedSdpaGraph(graphBuffer, size); - dense.has_value()) - { - executeRaggedAsDense(*dense, variantPack); - return; - } - } - executeAsIs(graphBuffer, size, variantPack); - } + const std::unordered_map& variantPack) override; bool requiresDeviceMemory() const override { @@ -62,325 +41,8 @@ class GpuReferenceGraphExecutor : public IReferenceGraphExecutor } private: - // Owns one device allocation; frees it after the device is idle. - struct DeviceBuffer - { - void* ptr = nullptr; - // Graph tensors have positive dimensions and empty sequence-length lists are not - // bound, so a zero-byte buffer is a caller bug. - explicit DeviceBuffer(size_t bytes) - { - if(bytes == 0) - { - throw std::invalid_argument("ragged SDPA densify: empty device buffer"); - } - hipdnn_gpu_ref::detail::throwOnHipError(hipMalloc(&ptr, bytes), - "ragged SDPA densify: hipMalloc"); - } - DeviceBuffer(const DeviceBuffer&) = delete; - DeviceBuffer& operator=(const DeviceBuffer&) = delete; - ~DeviceBuffer() - { - (void)hipDeviceSynchronize(); - (void)hipFree(ptr); - } - }; - - static void copy(void* to, const void* from, size_t bytes, hipMemcpyKind kind) - { - if(bytes > 0) - { - hipdnn_gpu_ref::detail::throwOnHipError(hipMemcpy(to, from, bytes, kind), - "ragged SDPA densify: hipMemcpy"); - } - } - - // The CPU executor's ragged-as-dense run with device buffers: offsets and packed inputs are - // read back, unpacked on the host and uploaded; dense outputs are read back, packed into a - // host copy of the packed output (so bytes outside live tokens are kept) and uploaded. - void executeRaggedAsDense(hipdnn_test_sdk::detail::DenseSdpaFromRagged& dense, - const std::unordered_map& variantPack) - { - namespace rd = hipdnn_test_sdk::detail; - const auto& primaries = dense.primaries; - std::vector> offsets; - std::vector> lengths; - for(const auto& primary : primaries) - { - std::vector table(static_cast(dense.batch + 1) - * primary.offsetElementSize); - copy(table.data(), - variantPack.at(primary.offsetUid), - table.size(), - hipMemcpyDeviceToHost); - offsets.push_back(rd::raggedElementOffsets(primary, table.data(), dense.batch)); - lengths.push_back(rd::raggedLengths(primary, offsets.back())); - } - const auto [qLengths, kvLengths] = rd::raggedGroupLengths(dense, lengths); - - auto densePack = variantPack; - std::vector> deviceBuffers; - std::vector> hostDense(primaries.size()); - for(size_t i = 0; i < primaries.size(); ++i) - { - hostDense[i].assign(rd::raggedDenseBytes(primaries[i]), 0); - if(!primaries[i].output) - { - std::vector packed( - rd::raggedPackedBytes(primaries[i], offsets[i], lengths[i])); - copy(packed.data(), - variantPack.at(primaries[i].uid), - packed.size(), - hipMemcpyDeviceToHost); - rd::copyRaggedTokens( - primaries[i], offsets[i], lengths[i], packed.data(), hostDense[i].data(), true); - } - deviceBuffers.push_back(std::make_unique(hostDense[i].size())); - copy(deviceBuffers.back()->ptr, - hostDense[i].data(), - hostDense[i].size(), - hipMemcpyHostToDevice); - densePack[primaries[i].uid] = deviceBuffers.back()->ptr; - } - const auto bindSeqLens = [&](std::optional added, - int64_t uid, - const std::vector& packed, - const char* name) { - if(packed.empty()) - { - return; - } - std::vector values(packed.size()); - if(added.has_value()) - { - values.assign(packed.begin(), packed.end()); - deviceBuffers.push_back( - std::make_unique(values.size() * sizeof(int32_t))); - copy(deviceBuffers.back()->ptr, - values.data(), - values.size() * sizeof(int32_t), - hipMemcpyHostToDevice); - densePack[uid] = deviceBuffers.back()->ptr; - } - else - { - copy(values.data(), - variantPack.at(uid), - values.size() * sizeof(int32_t), - hipMemcpyDeviceToHost); - rd::requireSeqLensWithinPacked(values.data(), packed, name); - } - }; - bindSeqLens(dense.addedSeqLenQUid, dense.seqLenQUid, qLengths, "seq_len_q"); - bindSeqLens(dense.addedSeqLenKvUid, dense.seqLenKvUid, kvLengths, "seq_len_kv"); - - executeAsIs(dense.graph.data(), dense.graph.size(), densePack); - - for(size_t i = 0; i < primaries.size(); ++i) - { - if(!primaries[i].output) - { - continue; - } - copy(hostDense[i].data(), - densePack.at(primaries[i].uid), - hostDense[i].size(), - hipMemcpyDeviceToHost); - std::vector packed( - rd::raggedPackedBytes(primaries[i], offsets[i], lengths[i])); - copy(packed.data(), - variantPack.at(primaries[i].uid), - packed.size(), - hipMemcpyDeviceToHost); - rd::copyRaggedTokens( - primaries[i], offsets[i], lengths[i], packed.data(), hostDense[i].data(), false); - copy(variantPack.at(primaries[i].uid), - packed.data(), - packed.size(), - hipMemcpyHostToDevice); - } - } - - bool isApplicableAsIs(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(!isNodeApplicable(node, graphWrap.getTensorMap())) - { - return false; - } - } - return true; - } - - void executeAsIs(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(buildPlanForNode(graphWrap, node)); - } - - std::vector> virtualTensors; - const auto variantPackWithVirtualTensors = populateVariantPackWithMissingVirtualTensors( - variantPack, graphWrap.getTensorMap(), virtualTensors); - - for(auto& executor : planExecutors) - { - executor->execute(variantPackWithVirtualTensors); - } - } - - 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 node with any ragged primary is keyed as ragged. A partly ragged node is declined - // by both plans (the ragged plan needs all four primaries ragged, and the dense plan - // rejects any ragged tensor); isApplicable then tries the dense rewrite. - const auto* sdpaAttributes = node.attributes_as_SdpaAttributes(); - if(sdpaAttributes != nullptr - && hipdnn_test_sdk::detail::hasRaggedSdpaPrimary(*sdpaAttributes, tensorMap)) - { - return detail::GpuSdpaRaggedFwdSignatureKey(node, tensorMap); - } - return detail::GpuSdpaFwdSignatureKey(node, tensorMap); - } - case NodeAttrs::SdpaBackwardAttributes: - return detail::GpuSdpaBwdSignatureKey(node, tensorMap); - - // Node types with no GPU plan yet - throw descriptive error - case NodeAttrs::BatchnormBackwardAttributes: - case NodeAttrs::ConvolutionBwdAttributes: - case NodeAttrs::ConvolutionWrwAttributes: - 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/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp b/dnn-providers/integration-tests/src/harness/gpu-graph-executor/detail/GpuPlanRegistrySignatureKey.hpp index 16669e978883..547af307cffb 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" @@ -53,23 +54,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/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/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/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/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 adbaf3408c3a..9461402fd0b1 100644 --- a/dnn-providers/integration-tests/tests/CMakeLists.txt +++ b/dnn-providers/integration-tests/tests/CMakeLists.txt @@ -3,14 +3,19 @@ # 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, 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 +# 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 @@ -21,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 @@ -64,6 +70,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/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/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 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/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/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 e0e7d37344a6..6056ba128b48 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 @@ -18,14 +21,15 @@ #include #include #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 +39,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 +356,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/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) diff --git a/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp b/dnn-providers/integration-tests/tests/TestVerificationPaths.cpp index 1b8dfc3dc7de..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 @@ -12,8 +13,8 @@ #include #include #include -#include +#include "harness/CpuReferenceGraphExecutorAdapter.hpp" #include "harness/bundle/BundleDiscovery.hpp" #include "harness/gpu-graph-executor/GpuReferenceGraphExecutor.hpp" @@ -91,10 +92,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 +118,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); 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 940aacbc9ea4..23af0b354bf7 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 @@ -79,22 +80,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); } };