[MIOpen] ROCM-24413 Fix CK dynamic-lib applicability tests

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

projects/miopen/test/gtest/gtest_common.hpp

@@ -34,6 +34,7 @@
 #include <tuple>
 #include <sstream>
 #include <string>
+#include <string_view>
 #include <vector>
 
 #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<Gpu, DevDescription>& GetAllKnownDevices();
 bool IsTestSupportedByDevice(Gpu supported_devs);
 

projects/miopen/test/gtest/unit_conv_solver.cpp

@@ -88,6 +88,27 @@ 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);
+}
+
+bool IsCKDynamicLibLoaded(std::string_view dev_name)
+{
+    return miopen::solver::CkImplLibLoader::Get(std::string{dev_name}).IsLoaded();
+}
+
 } // namespace
 
 //************************************************************************************
@@ -863,14 +884,18 @@ 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();
     }
     else if(params.check_xnack_disabled && get_handle_xnack())
     {
         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,
@@ -902,11 +927,22 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl(
 
     const auto problem = conv_config.GetProblemDescription(direction);
 
-    const auto all_known_devs = GetAllKnownDevices();
+    if(params.uses_ck_dynamic_lib)
+    {
+        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 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};
@@ -921,15 +957,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";

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;

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupBwdXdlops.cpp

@@ -50,8 +50,8 @@ auto GetDevApplicabilityConvCase()
 template <TestDataType type>
 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)
     {

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupFwdXdlops.cpp

@@ -61,8 +61,8 @@ auto GetDeterministicConvCase()
 template <TestDataType type>
 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)
     {

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemm3DGroupWrwXdlops.cpp

@@ -61,8 +61,8 @@ auto GetDeterministicConvCase()
 template <TestDataType type>
 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)
     {

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupBwdXdlops.cpp

@@ -61,8 +61,8 @@ auto GetDeterministicConvCase()
 template <TestDataType type>
 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)
     {

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupFwdXdlops.cpp

@@ -61,8 +61,8 @@ auto GetDeterministicConvCase()
 template <TestDataType type>
 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)
     {

projects/miopen/test/gtest/unit_conv_solver_ConvHipImplicitGemmGroupWrwXdlops.cpp

@@ -60,8 +60,8 @@ auto GetDeterministicConvCase()
 template <TestDataType type>
 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)
     {