From 2d5826abf48649d06258507b1789d90b888bfb18 Mon Sep 17 00:00:00 2001 From: Brian Harrison Date: Tue, 12 May 2026 13:25:05 -0400 Subject: [PATCH 1/2] [MIOpen] ROCM-24413 Fix CK dynamic-lib applicability tests --- projects/miopen/test/gtest/gtest_common.cpp | 61 +++++++++++-------- projects/miopen/test/gtest/gtest_common.hpp | 3 + .../miopen/test/gtest/unit_conv_solver.cpp | 55 ++++++++++++++--- .../unit_conv_solver_ConvCkGroupedConvFwd.cpp | 2 +- ...er_ConvHipImplicitGemm3DGroupBwdXdlops.cpp | 4 +- ...er_ConvHipImplicitGemm3DGroupFwdXdlops.cpp | 4 +- ...er_ConvHipImplicitGemm3DGroupWrwXdlops.cpp | 4 +- ...lver_ConvHipImplicitGemmGroupBwdXdlops.cpp | 4 +- ...lver_ConvHipImplicitGemmGroupFwdXdlops.cpp | 4 +- ...lver_ConvHipImplicitGemmGroupWrwXdlops.cpp | 4 +- 10 files changed, 95 insertions(+), 50 deletions(-) diff --git a/projects/miopen/test/gtest/gtest_common.cpp b/projects/miopen/test/gtest/gtest_common.cpp index 671d6a36c23c..cf8c2111d0e0 100644 --- a/projects/miopen/test/gtest/gtest_common.cpp +++ b/projects/miopen/test/gtest/gtest_common.cpp @@ -66,35 +66,42 @@ std::size_t MockHandle::GetMaxMemoryAllocSize() const bool MockHandle::CooperativeLaunchSupported() const { return false; } -Gpu GetDevGpuType() +std::string_view GetBaseDeviceName(std::string_view dev_name) { - const auto dev_name = get_handle().GetDeviceName(); - - static const auto dev = [&] { - if(dev_name == "gfx900") - return Gpu::gfx900; - else if(dev_name == "gfx906") - return Gpu::gfx906; - else if(dev_name == "gfx908") - return Gpu::gfx908; - else if(dev_name == "gfx90a") - return Gpu::gfx90A; - else if(dev_name == "gfx942") - return Gpu::gfx94X; - else if(miopen::StartsWith(dev_name, "gfx95")) - return Gpu::gfx950; - else if(miopen::StartsWith(dev_name, "gfx103")) - return Gpu::gfx103X; - else if(miopen::StartsWith(dev_name, "gfx110")) - return Gpu::gfx110X; - else if(miopen::StartsWith(dev_name, "gfx115")) - return Gpu::gfx115X; - else if(miopen::StartsWith(dev_name, "gfx120")) - return Gpu::gfx120X; - else - throw std::runtime_error("unknown_gpu"); - }(); + const auto suffix_pos = dev_name.find(':'); + return suffix_pos == std::string_view::npos ? dev_name : dev_name.substr(0, suffix_pos); +} +Gpu GetGpuType(const std::string& dev_name) +{ + const auto arch = std::string{GetBaseDeviceName(dev_name)}; + + if(arch == "gfx900") + return Gpu::gfx900; + if(arch == "gfx906") + return Gpu::gfx906; + if(arch == "gfx908") + return Gpu::gfx908; + if(arch == "gfx90a") + return Gpu::gfx90A; + if(arch == "gfx942") + return Gpu::gfx94X; + if(miopen::StartsWith(arch, "gfx95")) + return Gpu::gfx950; + if(miopen::StartsWith(arch, "gfx103")) + return Gpu::gfx103X; + if(miopen::StartsWith(arch, "gfx110")) + return Gpu::gfx110X; + if(miopen::StartsWith(arch, "gfx115")) + return Gpu::gfx115X; + if(miopen::StartsWith(arch, "gfx120")) + return Gpu::gfx120X; + throw std::runtime_error("unknown_gpu"); +} + +Gpu GetDevGpuType() +{ + static const auto dev = GetGpuType(get_handle().GetDeviceName()); return dev; } diff --git a/projects/miopen/test/gtest/gtest_common.hpp b/projects/miopen/test/gtest/gtest_common.hpp index 48051f36242e..dcbe36829790 100644 --- a/projects/miopen/test/gtest/gtest_common.hpp +++ b/projects/miopen/test/gtest/gtest_common.hpp @@ -34,6 +34,7 @@ #include #include #include +#include #include #include "../driver.hpp" @@ -230,6 +231,8 @@ class MockHandle final : public miopen::Handle }; Gpu GetDevGpuType(); +std::string_view GetBaseDeviceName(std::string_view dev_name); +Gpu GetGpuType(const std::string& dev_name); const std::multimap& GetAllKnownDevices(); bool IsTestSupportedByDevice(Gpu supported_devs); diff --git a/projects/miopen/test/gtest/unit_conv_solver.cpp b/projects/miopen/test/gtest/unit_conv_solver.cpp index 859b6f995d84..db5cdda8ade1 100644 --- a/projects/miopen/test/gtest/unit_conv_solver.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver.cpp @@ -88,6 +88,22 @@ bool IsDeviceSupported(Gpu supported_devs, Gpu dev) return false; } +bool IsDeviceExcluded(const UnitTestConvSolverParams& params, std::string_view dev_name) +{ + // Runtime device names can include target properties such as xnack/sramecc. + // Test exclusions are written against the bare architecture name. + return params.excluded_devices.count(GetBaseDeviceName(dev_name)) != 0; +} + +bool IsDeviceSupported(const UnitTestConvSolverParams& params, + Gpu dev, + std::string_view dev_name, + bool xnack_enabled) +{ + return IsDeviceSupported(params.supported_devs, dev) && !IsDeviceExcluded(params, dev_name) && + !(params.check_xnack_disabled && xnack_enabled); +} + } // namespace //************************************************************************************ @@ -863,7 +879,7 @@ void UnitTestConvSolverBase::SetUpImpl(const UnitTestConvSolverParams& params) { GTEST_SKIP(); } - else if(params.excluded_devices.count(get_handle().GetDeviceName()) > 0) + else if(IsDeviceExcluded(params, get_handle().GetDeviceName())) { GTEST_SKIP(); } @@ -902,6 +918,34 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl( const auto problem = conv_config.GetProblemDescription(direction); + // CK dynamic-library solvers can consult runtime-loaded plugins and HIP state. + // Mock-device contexts are not valid for that path, so check only the real device. + if(params.uses_ck_dynamic_lib) + { + auto&& handle = get_handle(); + const auto ctx = [&] { + auto tmp = miopen::ExecutionContext{&handle}; + problem.SetupFloats(tmp); + problem.SetupComputeType(tmp); + return tmp; + }(); + + const auto current_dev_name = ctx.GetStream().GetDeviceName(); + auto supported = IsDeviceSupported( + params, GetGpuType(current_dev_name), current_dev_name, get_handle_xnack()); + if(supported && !miopen::solver::CkImplLibLoader::Get(current_dev_name).IsLoaded()) + supported = false; + + const auto is_applicable = solver.IsApplicable(ctx, problem); + if(is_applicable != supported) + { + GTEST_FAIL() << current_dev_name << " is" << (is_applicable ? "" : " not") + << " applicable for " << solver.SolverDbId() << " but " + << (supported ? "" : "not ") << "marked as supported"; + } + return; + } + const auto all_known_devs = GetAllKnownDevices(); for(const auto& [dev, dev_descr] : all_known_devs) { @@ -921,15 +965,6 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl( // std::cout << "IsApplicable: " << is_applicable << std::endl; if(is_applicable != supported) { - // If the solver uses CK dynamic libraries and the library for this - // device wasn't built, the solver correctly reports not-applicable. - if(params.uses_ck_dynamic_lib && supported && !is_applicable) - { - const auto& loader = - miopen::solver::CkImplLibLoader::Get(std::string(dev_descr.name)); - if(!loader.IsLoaded()) - continue; - } GTEST_FAIL() << dev_descr << " is" << (is_applicable ? "" : " not") << " applicable for " << solver.SolverDbId() << " but " << (supported ? "" : "not ") << "marked as supported"; diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvCkGroupedConvFwd.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvCkGroupedConvFwd.cpp index 0a08dc66c667..cfd550eba0ba 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvCkGroupedConvFwd.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvCkGroupedConvFwd.cpp @@ -118,7 +118,7 @@ auto GetConvFullTestCases(miopenDataType_t datatype) auto GetTestParams(miopenDataType_t /*datatype*/) { // Solution requires 64-lane wavefronts and depends on the CK dynamic library -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +#if MIOPEN_BACKEND_HIP Gpu supportedDevices = Gpu::gfx908 | Gpu::gfx90A | Gpu::gfx94X | Gpu::gfx950; #else Gpu supportedDevices = Gpu::None; diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp index b8189aecc31e..5b88e623e53c 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp @@ -50,8 +50,8 @@ auto GetDevApplicabilityConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp index 9d4e6004ef7e..af65aa97244f 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp @@ -61,8 +61,8 @@ auto GetDeterministicConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp index 571b0e1c9384..beaf2f4ff27a 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp @@ -61,8 +61,8 @@ auto GetDeterministicConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp index 1eef20d26ee6..9e78a5d7d527 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp @@ -61,8 +61,8 @@ auto GetDeterministicConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp index 69af06af4687..21163c311dbc 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp @@ -61,8 +61,8 @@ auto GetDeterministicConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { diff --git a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp index bd7ef627c703..2c05f2ac08b1 100644 --- a/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp @@ -60,8 +60,8 @@ auto GetDeterministicConvCase() template miopen::unit_tests::UnitTestConvSolverParams GetTestParams() { -// If MIOpen is built without CK these tests will fail, skip them to avoid failing -#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL +// CK dynamic-library tests are HIP-only; runtime plugin availability is checked by the harness. +#if MIOPEN_BACKEND_HIP Gpu supportedDevices; if constexpr(type == TestDataType::FP32) { From c0490e8b5a6929cf8371f77f771e178a033c34d1 Mon Sep 17 00:00:00 2001 From: Brian Harrison Date: Tue, 12 May 2026 18:21:32 -0400 Subject: [PATCH 2/2] MIOpen: skip CK dynamic tests without device library --- .../miopen/test/gtest/unit_conv_solver.cpp | 46 ++++++++----------- 1 file changed, 19 insertions(+), 27 deletions(-) diff --git a/projects/miopen/test/gtest/unit_conv_solver.cpp b/projects/miopen/test/gtest/unit_conv_solver.cpp index db5cdda8ade1..70e54ca7c563 100644 --- a/projects/miopen/test/gtest/unit_conv_solver.cpp +++ b/projects/miopen/test/gtest/unit_conv_solver.cpp @@ -104,6 +104,11 @@ bool IsDeviceSupported(const UnitTestConvSolverParams& params, !(params.check_xnack_disabled && xnack_enabled); } +bool IsCKDynamicLibLoaded(std::string_view dev_name) +{ + return miopen::solver::CkImplLibLoader::Get(std::string{dev_name}).IsLoaded(); +} + } // namespace //************************************************************************************ @@ -887,6 +892,10 @@ void UnitTestConvSolverBase::SetUpImpl(const UnitTestConvSolverParams& params) { GTEST_SKIP(); } + else if(params.uses_ck_dynamic_lib && !IsCKDynamicLibLoaded(get_handle().GetDeviceName())) + { + GTEST_SKIP() << "CK dynamic library is not available for " << get_handle().GetDeviceName(); + } } void UnitTestConvSolverBase::RunTestImpl(const miopen::solver::conv::ConvSolverInterface& solver, @@ -918,39 +927,22 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl( const auto problem = conv_config.GetProblemDescription(direction); - // CK dynamic-library solvers can consult runtime-loaded plugins and HIP state. - // Mock-device contexts are not valid for that path, so check only the real device. if(params.uses_ck_dynamic_lib) { - auto&& handle = get_handle(); - const auto ctx = [&] { - auto tmp = miopen::ExecutionContext{&handle}; - problem.SetupFloats(tmp); - problem.SetupComputeType(tmp); - return tmp; - }(); - - const auto current_dev_name = ctx.GetStream().GetDeviceName(); - auto supported = IsDeviceSupported( - params, GetGpuType(current_dev_name), current_dev_name, get_handle_xnack()); - if(supported && !miopen::solver::CkImplLibLoader::Get(current_dev_name).IsLoaded()) - supported = false; - - const auto is_applicable = solver.IsApplicable(ctx, problem); - if(is_applicable != supported) - { - GTEST_FAIL() << current_dev_name << " is" << (is_applicable ? "" : " not") - << " applicable for " << solver.SolverDbId() << " but " - << (supported ? "" : "not ") << "marked as supported"; - } - return; + const auto current_dev_name = get_handle().GetDeviceName(); + if(!IsCKDynamicLibLoaded(current_dev_name)) + GTEST_SKIP() << "CK dynamic library is not available for " << current_dev_name; } - const auto all_known_devs = GetAllKnownDevices(); + const auto current_dev_name = get_handle().GetDeviceName(); + const auto current_dev_base_name = GetBaseDeviceName(current_dev_name); + const auto all_known_devs = GetAllKnownDevices(); for(const auto& [dev, dev_descr] : all_known_devs) { - const auto supported = IsDeviceSupported(params.supported_devs, dev) && - params.excluded_devices.count(dev_descr.name) == 0; + if(params.uses_ck_dynamic_lib && current_dev_base_name != dev_descr.name) + continue; + + const auto supported = IsDeviceSupported(params, dev, dev_descr.name, false); // std::cout << "Test " << dev_descr << " (supported: " << supported << ")" << std::endl; auto handle = MockHandle{dev_descr, params.check_xnack_disabled};