perf(hipdnn): speed up superbuild CI with parallel ctest and compiled reference executors

.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.

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})

dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationHarness.hpp

@@ -13,9 +13,10 @@
 #include <hipdnn_test_sdk/utilities/CpuFpReferenceValidation.hpp>
 #include <hipdnn_test_sdk/utilities/SdkFrontendTypeConversions.hpp>
 #include <hipdnn_test_sdk/utilities/TestUtilities.hpp>
-#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceGraphExecutor.hpp>
 #include <hipdnn_test_sdk/utilities/cpu_graph_executor/GraphTensorBundle.hpp>
 
+#include <cstddef>
+#include <cstdint>
 #include <limits>
 #include <memory>
 #include <string>
@@ -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<int64_t, void*>& variantPack);
+
 // NOLINTBEGIN (portability-template-virtual-member-function)
 template <typename DataType, typename TestCaseType>
 class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<TestCaseType>
@@ -412,7 +420,7 @@ class IntegrationGraphVerificationHarness : public ::testing::TestWithParam<Test
         auto [serializedGraph, serErr] = graph.to_binary();
         ASSERT_TRUE(serErr.is_good()) << serErr.get_message();
 
-        hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute(
+        executeCpuReferenceGraph(
             serializedGraph.data(), serializedGraph.size(), bundle.toHostVariantPack());
     }
 

dnn-providers/hip-kernel-provider/src/integration_tests/IntegrationGraphVerificationReference.cpp

@@ -0,0 +1,18 @@
+// Copyright © Advanced Micro Devices, Inc., or its affiliates.
+// SPDX-License-Identifier: MIT
+
+#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceGraphExecutor.hpp>
+
+#include "IntegrationGraphVerificationHarness.hpp"
+
+namespace hip_kernel_provider::test_utilities
+{
+
+void executeCpuReferenceGraph(void* graphBuffer,
+                              size_t size,
+                              const std::unordered_map<int64_t, void*>& variantPack)
+{
+    hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor().execute(graphBuffer, size, variantPack);
+}
+
+} // namespace hip_kernel_provider::test_utilities

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
-        $<$<BOOL:${USE_ROCRAND}>:roc::rocrand>
-        $<$<PLATFORM_ID:Windows>: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
+        $<$<BOOL:${USE_ROCRAND}>:roc::rocrand>
+        $<$<PLATFORM_ID:Windows>: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)
 

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 <string>
+
+#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceGraphExecutor.hpp>
+#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceNotApplicableError.hpp>
+
+namespace hipdnn_integration_tests
+{
+
+struct CpuReferenceGraphExecutorAdapter::Impl
+{
+    hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor executor;
+};
+
+CpuReferenceGraphExecutorAdapter::CpuReferenceGraphExecutorAdapter()
+    : _impl(std::make_unique<Impl>())
+{
+}
+
+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<int64_t, void*>& 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

dnn-providers/integration-tests/src/harness/CpuReferenceGraphExecutorAdapter.hpp

@@ -3,55 +3,48 @@
 
 #pragma once
 
-#include <stdexcept>
-#include <string>
-
-#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceGraphExecutor.hpp>
-#include <hipdnn_test_sdk/utilities/cpu_graph_executor/CpuReferenceNotApplicableError.hpp>
+#include <cstddef>
+#include <cstdint>
+#include <memory>
+#include <unordered_map>
 
 #include "IReferenceGraphExecutor.hpp"
 #include "ReferenceCapabilityError.hpp"
 
 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<int64_t, void*>& 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<int64_t, void*>& variantPack) override;
 
     bool requiresDeviceMemory() const override
     {
         return false;
     }
 
 private:
-    hipdnn_test_sdk::utilities::CpuReferenceGraphExecutor _executor;
+    struct Impl;
+    std::unique_ptr<Impl> _impl;
 };
 
 } // namespace hipdnn_integration_tests