diff --git a/src/hip/handlehip.cpp b/src/hip/handlehip.cpp index 0a033900e4..9a21a6dc9f 100644 --- a/src/hip/handlehip.cpp +++ b/src/hip/handlehip.cpp @@ -782,7 +782,7 @@ bool Handle::CooperativeLaunchSupported() const std::string Handle::GetDeviceNameImpl() const { return this->impl->get_device_name(); } -std::string Handle::GetDeviceName() const { return this->impl->target_properties.Name(); } +std::string Handle::GetDeviceName() const { return this->GetTargetProperties().Name(); } const TargetProperties& Handle::GetTargetProperties() const { diff --git a/src/include/miopen/conv/solvers.hpp b/src/include/miopen/conv/solvers.hpp index 4890519a00..bcb39b6d72 100644 --- a/src/include/miopen/conv/solvers.hpp +++ b/src/include/miopen/conv/solvers.hpp @@ -331,37 +331,37 @@ struct ConvAsm1x1UV2 final : ConvTunableSolver const PerformanceConfigConvAsm1x1UV2&) const override; }; -struct ConvAsm5x10u2v2f1 final : ConvSolver +struct MIOPEN_INTERNALS_EXPORT ConvAsm5x10u2v2f1 final : ConvSolver { const std::string& SolverDbId() const override { return GetSolverDbId(); } - MIOPEN_INTERNALS_EXPORT bool - IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; - MIOPEN_INTERNALS_EXPORT ConvSolution - GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; + bool IsApplicable(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; + ConvSolution GetSolution(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; }; -struct ConvAsm5x10u2v2b1 final : ConvSolver +struct MIOPEN_INTERNALS_EXPORT ConvAsm5x10u2v2b1 final : ConvSolver { const std::string& SolverDbId() const override { return GetSolverDbId(); } - MIOPEN_INTERNALS_EXPORT bool - IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; - MIOPEN_INTERNALS_EXPORT ConvSolution - GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; + bool IsApplicable(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; + ConvSolution GetSolution(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; }; -struct ConvAsm7x7c3h224w224k64u2v2p3q3f1 final : ConvSolver +struct MIOPEN_INTERNALS_EXPORT ConvAsm7x7c3h224w224k64u2v2p3q3f1 final : ConvSolver { const std::string& SolverDbId() const override { return GetSolverDbId(); } - MIOPEN_INTERNALS_EXPORT bool - IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; - MIOPEN_INTERNALS_EXPORT ConvSolution - GetSolution(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override; + bool IsApplicable(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; + ConvSolution GetSolution(const ExecutionContext&, + const miopen::conv::ProblemDescription&) const override; }; struct ConvOclDirectFwd11x11 final : ConvSolver diff --git a/src/include/miopen/fusion/utils.hpp b/src/include/miopen/fusion/utils.hpp index 38e1d69789..3c44dfe7ae 100644 --- a/src/include/miopen/fusion/utils.hpp +++ b/src/include/miopen/fusion/utils.hpp @@ -92,7 +92,7 @@ inline bool WinoCommonIsApplicable(const FusionContext& context, const FusionDes return false; if(!conv_problem.IsDirectionForward()) return false; - const auto target = conv_ctx.GetStream().GetTargetProperties(); + const auto& target = conv_ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/include/miopen/handle.hpp b/src/include/miopen/handle.hpp index 31096d0a34..b12b267bf5 100644 --- a/src/include/miopen/handle.hpp +++ b/src/include/miopen/handle.hpp @@ -177,8 +177,8 @@ struct MIOPEN_EXPORT Handle : miopenHandle virtual std::size_t GetMaxMemoryAllocSize() const; virtual bool CooperativeLaunchSupported() const; - virtual std::string GetDeviceName() const; - const TargetProperties& GetTargetProperties() const; + std::string GetDeviceName() const; + virtual const TargetProperties& GetTargetProperties() const; private: std::string GetDeviceNameImpl() const; diff --git a/src/include/miopen/target_properties.hpp b/src/include/miopen/target_properties.hpp index a66ebc27be..2b8b04b413 100644 --- a/src/include/miopen/target_properties.hpp +++ b/src/include/miopen/target_properties.hpp @@ -37,9 +37,11 @@ struct Handle; struct TargetProperties { - const std::string& Name() const { return name; } + virtual ~TargetProperties() = default; + + virtual const std::string& Name() const { return name; } const std::string& DbId() const { return dbId; } - boost::optional Xnack() const { return xnack; } + virtual boost::optional Xnack() const { return xnack; } boost::optional Sramecc() const { return sramecc; } boost::optional SrameccReported() const { return sramecc_reported; } static std::size_t GetMaxWaveScratchSize() { return MaxWaveScratchSize; } diff --git a/src/nogpu/handle.cpp b/src/nogpu/handle.cpp index fbe292ef50..b669632c31 100644 --- a/src/nogpu/handle.cpp +++ b/src/nogpu/handle.cpp @@ -274,7 +274,7 @@ const TargetProperties& Handle::GetTargetProperties() const } std::string Handle::GetDeviceNameImpl() const { return this->impl->device_name; } -std::string Handle::GetDeviceName() const { return this->impl->target_properties.Name(); } +std::string Handle::GetDeviceName() const { return this->GetTargetProperties().Name(); } std::ostream& Handle::Print(std::ostream& os) const { diff --git a/src/ocl/handleocl.cpp b/src/ocl/handleocl.cpp index 425f3a0aee..2ef8ef1014 100644 --- a/src/ocl/handleocl.cpp +++ b/src/ocl/handleocl.cpp @@ -469,7 +469,7 @@ std::size_t Handle::GetGlobalMemorySize() const std::string Handle::GetDeviceNameImpl() const { return this->impl->get_device_name(); } -std::string Handle::GetDeviceName() const { return this->impl->target_properties.Name(); } +std::string Handle::GetDeviceName() const { return this->GetTargetProperties().Name(); } const TargetProperties& Handle::GetTargetProperties() const { diff --git a/src/solver/conv/conv_MP_bidirectional_winograd.cpp b/src/solver/conv/conv_MP_bidirectional_winograd.cpp index 1353652289..983dbaa744 100644 --- a/src/solver/conv/conv_MP_bidirectional_winograd.cpp +++ b/src/solver/conv/conv_MP_bidirectional_winograd.cpp @@ -193,7 +193,7 @@ static bool IsApplicableTransform(const ExecutionContext& ctx, const ProblemDesc if(!(problem.IsFp32() || problem.IsFp16())) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_1x1u.cpp b/src/solver/conv/conv_asm_1x1u.cpp index ddb7518057..587145c627 100644 --- a/src/solver/conv/conv_asm_1x1u.cpp +++ b/src/solver/conv/conv_asm_1x1u.cpp @@ -545,7 +545,7 @@ bool ConvAsm1x1U::IsApplicable(const ExecutionContext& ctx, const ProblemDescrip if(problem.IsTensorsCasted()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_1x1u_stride2.cpp b/src/solver/conv/conv_asm_1x1u_stride2.cpp index 2537aca6cf..937ab19151 100644 --- a/src/solver/conv/conv_asm_1x1u_stride2.cpp +++ b/src/solver/conv/conv_asm_1x1u_stride2.cpp @@ -506,7 +506,7 @@ bool ConvAsm1x1UV2::IsApplicable(const ExecutionContext& ctx, if(problem.IsTensorsCasted()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_3x3u.cpp b/src/solver/conv/conv_asm_3x3u.cpp index e77bd621eb..839517e3c7 100644 --- a/src/solver/conv/conv_asm_3x3u.cpp +++ b/src/solver/conv/conv_asm_3x3u.cpp @@ -192,7 +192,7 @@ bool ConvAsm3x3U::IsApplicable(const ExecutionContext& ctx, const ProblemDescrip if(!ctx.rmv.IsV2orV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_5x10u2v2b1.cpp b/src/solver/conv/conv_asm_5x10u2v2b1.cpp index 47438889ca..61c86b92af 100644 --- a/src/solver/conv/conv_asm_5x10u2v2b1.cpp +++ b/src/solver/conv/conv_asm_5x10u2v2b1.cpp @@ -58,7 +58,7 @@ bool ConvAsm5x10u2v2b1::IsApplicable(const ExecutionContext& ctx, if(!ctx.rmv.IsV2orV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_5x10u2v2f1.cpp b/src/solver/conv/conv_asm_5x10u2v2f1.cpp index 3b37f2f316..63a72e4bc4 100644 --- a/src/solver/conv/conv_asm_5x10u2v2f1.cpp +++ b/src/solver/conv/conv_asm_5x10u2v2f1.cpp @@ -58,7 +58,7 @@ bool ConvAsm5x10u2v2f1::IsApplicable(const ExecutionContext& ctx, if(!ctx.rmv.IsV2orV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp b/src/solver/conv/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp index 628f62e6da..5e5a0d0a8c 100644 --- a/src/solver/conv/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp +++ b/src/solver/conv/conv_asm_7x7c3h224w224k64u2v2p3q3f1.cpp @@ -63,7 +63,7 @@ bool ConvAsm7x7c3h224w224k64u2v2p3q3f1::IsApplicable(const ExecutionContext& ctx if(problem.IsTensorsCasted()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp b/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp index 3747d172c7..5348656009 100644 --- a/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp +++ b/src/solver/conv/conv_asm_dir_BwdWrW1x1.cpp @@ -496,7 +496,7 @@ bool ConvAsmBwdWrW1x1::IsApplicable(const ExecutionContext& ctx, if(problem.IsTensorsCasted()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp b/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp index 7ac59dc686..44cfa03527 100644 --- a/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp +++ b/src/solver/conv/conv_asm_dir_BwdWrW3x3.cpp @@ -410,7 +410,7 @@ bool ConvAsmBwdWrW3x3::IsApplicable(const ExecutionContext& ctx, if(!ctx.rmv.IsV2orV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp b/src/solver/conv/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp index d75ca4f2a8..8657b161d3 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_bwd_v4r1_dynamic.cpp @@ -170,7 +170,7 @@ bool ConvAsmImplicitGemmV4R1DynamicBwd::IsApplicable(const ExecutionContext& ctx if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd.cpp index b1c737a195..bb27d59bf5 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd.cpp @@ -1017,7 +1017,7 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlops::IsApplicable(const ExecutionContext if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp index e20b4681b6..e5bd7094b7 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_bwd_nhwc.cpp @@ -969,7 +969,7 @@ bool ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC::IsApplicable( if(!ctx.rmv.IsV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; // NOLINT (readability-simplify-boolean-expr) diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd.cpp index e6363b11ed..0e5f77369b 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd.cpp @@ -1550,7 +1550,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlops::IsApplicable(const ExecutionContext } #endif - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp index ecfe686fe2..4d2c80923d 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nchwc.cpp @@ -585,7 +585,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC::IsApplicable( if(!ctx.rmv.IsV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; // NOLINT (readability-simplify-boolean-expr) diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp index f821e118e4..fc8848c02a 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_fwd_nhwc.cpp @@ -912,7 +912,7 @@ bool ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC::IsApplicable( if(!ctx.rmv.IsV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; // NOLINT (readability-simplify-boolean-expr) diff --git a/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index fac718bb61..717b4f60a8 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -918,7 +918,7 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::IsApplicable( if(!ctx.rmv.IsV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; // NOLINT (readability-simplify-boolean-expr) diff --git a/src/solver/conv/conv_asm_implicit_gemm_v4r1_dynamic.cpp b/src/solver/conv/conv_asm_implicit_gemm_v4r1_dynamic.cpp index e2ece291b0..dcb8c81527 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_v4r1_dynamic.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_v4r1_dynamic.cpp @@ -320,7 +320,7 @@ bool ConvAsmImplicitGemmV4R1DynamicFwd::IsApplicable(const ExecutionContext& ctx if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; auto tunables = GetImplicitGemmV4R1DynamicTunables(); @@ -366,7 +366,7 @@ bool ConvAsmImplicitGemmV4R1DynamicFwd_1x1::IsApplicable(const ExecutionContext& if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; auto tunables = GetImplicitGemmV4R1DynamicTunables(); diff --git a/src/solver/conv/conv_asm_implicit_gemm_wrw_gtc_dynamic_xdlops.cpp b/src/solver/conv/conv_asm_implicit_gemm_wrw_gtc_dynamic_xdlops.cpp index 240afb4132..cb377385c2 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_wrw_gtc_dynamic_xdlops.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_wrw_gtc_dynamic_xdlops.cpp @@ -861,7 +861,7 @@ bool ConvAsmImplicitGemmGTCDynamicWrwXdlops::IsApplicable(const ExecutionContext if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; bool is_valid; diff --git a/src/solver/conv/conv_asm_implicit_gemm_wrw_v4r1_dynamic.cpp b/src/solver/conv/conv_asm_implicit_gemm_wrw_v4r1_dynamic.cpp index d9907eea68..0a48b78c66 100644 --- a/src/solver/conv/conv_asm_implicit_gemm_wrw_v4r1_dynamic.cpp +++ b/src/solver/conv/conv_asm_implicit_gemm_wrw_v4r1_dynamic.cpp @@ -341,7 +341,7 @@ bool ConvAsmImplicitGemmV4R1DynamicWrw::IsApplicable(const ExecutionContext& ctx if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; std::string kernel_name; diff --git a/src/solver/conv/conv_bin_wino3x3U.cpp b/src/solver/conv/conv_bin_wino3x3U.cpp index 651c3a23dd..daf824a4ac 100644 --- a/src/solver/conv/conv_bin_wino3x3U.cpp +++ b/src/solver/conv/conv_bin_wino3x3U.cpp @@ -55,7 +55,7 @@ bool ConvBinWinograd3x3U::IsApplicable(const ExecutionContext& ctx, if(!(ctx.rmv.IsV2orV3() && ctx.use_asm_kernels)) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_bin_winoRxS.cpp b/src/solver/conv/conv_bin_winoRxS.cpp index ef97fe1659..17586cdc3e 100644 --- a/src/solver/conv/conv_bin_winoRxS.cpp +++ b/src/solver/conv/conv_bin_winoRxS.cpp @@ -249,7 +249,7 @@ bool ConvBinWinogradRxS::IsApplicable(const ExecutionContext& ctx, if(!ctx.rmv.IsV2orV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_multipass_wino3x3WrW.cpp b/src/solver/conv/conv_multipass_wino3x3WrW.cpp index f6739b9da5..a9d9a5f052 100644 --- a/src/solver/conv/conv_multipass_wino3x3WrW.cpp +++ b/src/solver/conv/conv_multipass_wino3x3WrW.cpp @@ -477,7 +477,7 @@ bool ConvWinograd3x3MultipassWrW if(!problem.IsLayoutDefault()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/src/solver/conv/conv_winoRxS.cpp b/src/solver/conv/conv_winoRxS.cpp index 9a44e8590e..3e6ac5c5b0 100644 --- a/src/solver/conv/conv_winoRxS.cpp +++ b/src/solver/conv/conv_winoRxS.cpp @@ -671,7 +671,7 @@ static bool IsApplicableBase(const ExecutionContext& ctx, const ProblemDescripti if(!ctx.rmv.IsV3()) return false; - const auto target = ctx.GetStream().GetTargetProperties(); + const auto& target = ctx.GetStream().GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/test/gtest/conv_igemm_dynamic_xdlops.cpp b/test/gtest/conv_igemm_dynamic_xdlops.cpp index 60a991767c..e75d04a310 100644 --- a/test/gtest/conv_igemm_dynamic_xdlops.cpp +++ b/test/gtest/conv_igemm_dynamic_xdlops.cpp @@ -105,7 +105,7 @@ using TestCase = decltype(GetTestCases())::value_type; bool IsTestSupportedForDevice(const miopen::Handle& handle) { - const auto target = handle.GetTargetProperties(); + const auto& target = handle.GetTargetProperties(); if(target.Xnack() && *target.Xnack()) return false; using e_mask = enabled; diff --git a/test/gtest/conv_igemm_dynamic_xdlops_half.cpp b/test/gtest/conv_igemm_dynamic_xdlops_half.cpp index f737db9ae5..1eb672831d 100644 --- a/test/gtest/conv_igemm_dynamic_xdlops_half.cpp +++ b/test/gtest/conv_igemm_dynamic_xdlops_half.cpp @@ -99,7 +99,7 @@ void Run2dDriver(miopenDataType_t prec) bool IsTestSupportedForDevice(const miopen::Handle& handle) { - const auto target = handle.GetTargetProperties(); + const auto& target = handle.GetTargetProperties(); std::string devName = handle.GetDeviceName(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp index ae6011604b..93c702d3a6 100644 --- a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp +++ b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp @@ -100,7 +100,7 @@ void Run2dDriver(miopenDataType_t prec) bool IsTestSupportedForDevice(const miopen::Handle& handle) { - const auto target = handle.GetTargetProperties(); + const auto& target = handle.GetTargetProperties(); std::string devName = handle.GetDeviceName(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp index 03907f773e..f578e61346 100644 --- a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp +++ b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp @@ -111,7 +111,7 @@ void Run2dDriver(miopenDataType_t prec) bool IsTestSupportedForDevice(const miopen::Handle& handle) { - const auto target = handle.GetTargetProperties(); + const auto& target = handle.GetTargetProperties(); std::string devName = handle.GetDeviceName(); if(target.Xnack() && *target.Xnack()) return false; diff --git a/test/gtest/gtest_common.cpp b/test/gtest/gtest_common.cpp index b058c7a165..ceaada3cda 100644 --- a/test/gtest/gtest_common.cpp +++ b/test/gtest/gtest_common.cpp @@ -31,9 +31,30 @@ std::ostream& operator<<(std::ostream& os, const DevDescription& dd) return os << dd.name << "(" << dd.cu_cnt << ")"; } -MockHandle::MockHandle(const DevDescription& dev_description) : dev_descr{dev_description} {} +MockTargetProperties::MockTargetProperties(const TargetProperties& target_properties, + const DevDescription& dev_description, + bool disable_xnack) + : TargetProperties{target_properties}, name{dev_description.name}, xnack_disabled{disable_xnack} +{ +} + +const std::string& MockTargetProperties::Name() const { return name; } + +boost::optional MockTargetProperties::Xnack() const +{ + return xnack_disabled ? boost::none : TargetProperties::Xnack(); +} -std::string MockHandle::GetDeviceName() const { return std::string{dev_descr.name}; } +MockHandle::MockHandle(const DevDescription& dev_description, bool disable_xnack) + : dev_descr{dev_description}, + target_properties{Handle::GetTargetProperties(), dev_description, disable_xnack} +{ +} + +const miopen::TargetProperties& MockHandle::GetTargetProperties() const +{ + return target_properties; +} std::size_t MockHandle::GetMaxComputeUnits() const { return dev_descr.cu_cnt; } diff --git a/test/gtest/gtest_common.hpp b/test/gtest/gtest_common.hpp index cc414475e3..220bd4226b 100644 --- a/test/gtest/gtest_common.hpp +++ b/test/gtest/gtest_common.hpp @@ -105,19 +105,36 @@ struct DevDescription friend std::ostream& operator<<(std::ostream& os, const DevDescription& dd); }; +class MockTargetProperties final : public miopen::TargetProperties +{ +public: + MockTargetProperties(const TargetProperties& target_properties, + const DevDescription& dev_description, + bool disable_xnack); + + // Add additional methods here if needed + const std::string& Name() const override; + boost::optional Xnack() const override; + +private: + std::string name; + bool xnack_disabled; +}; + class MockHandle final : public miopen::Handle { public: - MockHandle(const DevDescription& dev_description); + MockHandle(const DevDescription& dev_description, bool disable_xnack); // Add additional methods here if needed - std::string GetDeviceName() const override; + const miopen::TargetProperties& GetTargetProperties() const override; std::size_t GetMaxComputeUnits() const override; std::size_t GetMaxMemoryAllocSize() const override; bool CooperativeLaunchSupported() const override; private: DevDescription dev_descr; + MockTargetProperties target_properties; }; Gpu GetDevGpuType(); diff --git a/test/gtest/smoke_solver_ConvAsm_5x10_7x7.cpp b/test/gtest/smoke_solver_ConvAsm_5x10_7x7.cpp deleted file mode 100644 index 1d2d0dbba6..0000000000 --- a/test/gtest/smoke_solver_ConvAsm_5x10_7x7.cpp +++ /dev/null @@ -1,89 +0,0 @@ -/******************************************************************************* - * - * MIT License - * - * Copyright (c) 2023 Advanced Micro Devices, Inc. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in all - * copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE - * SOFTWARE. - * - *******************************************************************************/ -#include -#include - -#include "gtest_common.hpp" - -#include "../conv2d.hpp" - -namespace { - -auto GetTestCases() -{ - const auto env5x10f = std::tuple{std::pair{MIOPEN_FIND_MODE, "normal"}, - std::pair{MIOPEN_DEBUG_FIND_ONLY_SOLVER, "ConvAsm5x10u2v2f1"}}; - const auto env5x10b = std::tuple{std::pair{MIOPEN_FIND_MODE, "normal"}, - std::pair{MIOPEN_DEBUG_FIND_ONLY_SOLVER, "ConvAsm5x10u2v2b1"}}; - const auto env7x7 = - std::tuple{std::pair{MIOPEN_FIND_MODE, "normal"}, - std::pair{MIOPEN_DEBUG_FIND_ONLY_SOLVER, "ConvAsm7x7c3h224w224k64u2v2p3q3f1"}}; - - const std::string vf = " --verbose --disable-backward-data --disable-backward-weights"; - const std::string vb = " --verbose --disable-forward --disable-backward-weights"; - - return std::vector{ - // clang-format off - std::pair{env5x10f, vf + " --input 1 1 5 10 --weights 16 1 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, - std::pair{env5x10b, vb + " --input 1 1 16 160 --weights 16 16 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, - std::pair{env7x7, vf + " --input 1 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1"} - // clang-format on - }; -} - -using TestCase = decltype(GetTestCases())::value_type; - -bool SkipTest() { return get_handle_xnack(); } - -bool IsTestSupportedForDevice() -{ - // GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 - using e_mask = enabled; - using d_mask = disabled; - return ::IsTestSupportedForDevMask(); -} - -} // namespace - -class GPU_Conv2dDefaultAsm_5x10_7x7_FP32 : public FloatTestCase> -{ -}; - -TEST_P(GPU_Conv2dDefaultAsm_5x10_7x7_FP32, FloatTest_smoke_solver_ConvAsm_5x10_7x7) -{ - if(IsTestSupportedForDevice() && !SkipTest()) - { - invoke_with_params(default_check); - } - else - { - GTEST_SKIP(); - } -}; - -INSTANTIATE_TEST_SUITE_P(Smoke, - GPU_Conv2dDefaultAsm_5x10_7x7_FP32, - testing::Values(GetTestCases())); diff --git a/test/gtest/unit_conv_solver.cpp b/test/gtest/unit_conv_solver.cpp index fecf48cc1a..acc492c077 100644 --- a/test/gtest/unit_conv_solver.cpp +++ b/test/gtest/unit_conv_solver.cpp @@ -768,7 +768,7 @@ void UnitTestConvSolverDevApplicabilityBase::RunTestImpl( const auto supported = IsDeviceSupported(params.supported_devs, dev); // std::cout << "Test " << dev_descr << " (supported: " << supported << ")" << std::endl; - auto handle = MockHandle{dev_descr}; + auto handle = MockHandle{dev_descr, params.check_xnack_disabled}; const auto ctx = [&] { auto tmp = miopen::ExecutionContext{&handle}; problem.SetupFloats(tmp); diff --git a/test/gtest/unit_conv_solver_ConvAsm5x10u2v2b1.cpp b/test/gtest/unit_conv_solver_ConvAsm5x10u2v2b1.cpp new file mode 100644 index 0000000000..23e5e610b6 --- /dev/null +++ b/test/gtest/unit_conv_solver_ConvAsm5x10u2v2b1.cpp @@ -0,0 +1,82 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2025 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ + +#include "unit_conv_solver.hpp" + +namespace { + +auto GetConvTestCases(miopenDataType_t datatype) +{ + using TestCase = miopen::unit_tests::ConvTestCase; + + return std::vector{ + // clang-format off + TestCase{{1, 16, 36, 284}, {16, 16, 5, 10}, {0, 0}, {2, 2}, {1, 1}, datatype}, + // clang-format on + }; +} + +const auto& GetTestParams() +{ + static const auto params = [] { + // gfx90A is not enabled because of WORKAROUND_ISSUE_1146 + Gpu supported_gpus = Gpu::gfx900 | Gpu::gfx906 | Gpu::gfx908; + auto p = miopen::unit_tests::UnitTestConvSolverParams(supported_gpus); + p.CheckXnackDisabled(); + return p; + }(); + return params; +} + +} // namespace + +using GPU_UnitTestConvSolverAsm5x10u2v2b1Bwd_FP32 = GPU_UnitTestConvSolverBwd_FP32; + +using CPU_UnitTestConvSolverAsm5x10u2v2b1DevApplicabilityBwd_NONE = + CPU_UnitTestConvSolverDevApplicabilityBwd_NONE; + +TEST_P(GPU_UnitTestConvSolverAsm5x10u2v2b1Bwd_FP32, ConvAsm5x10u2v2b1) +{ + this->RunTest(miopen::solver::conv::ConvAsm5x10u2v2b1{}); +}; + +TEST_P(CPU_UnitTestConvSolverAsm5x10u2v2b1DevApplicabilityBwd_NONE, ConvAsm5x10u2v2b1) +{ + this->RunTest(miopen::solver::conv::ConvAsm5x10u2v2b1{}); +}; + +// Smoke tests +INSTANTIATE_TEST_SUITE_P(Smoke, + GPU_UnitTestConvSolverAsm5x10u2v2b1Bwd_FP32, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(miopenConvolutionAlgoDirect), + testing::ValuesIn(GetConvTestCases(miopenFloat)))); + +// Device applicability test +INSTANTIATE_TEST_SUITE_P(Smoke, + CPU_UnitTestConvSolverAsm5x10u2v2b1DevApplicabilityBwd_NONE, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(GetConvTestCases(miopenFloat)[0]))); diff --git a/test/gtest/unit_conv_solver_ConvAsm5x10u2v2f1.cpp b/test/gtest/unit_conv_solver_ConvAsm5x10u2v2f1.cpp new file mode 100644 index 0000000000..defea5d6c6 --- /dev/null +++ b/test/gtest/unit_conv_solver_ConvAsm5x10u2v2f1.cpp @@ -0,0 +1,82 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2025 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ + +#include "unit_conv_solver.hpp" + +namespace { + +auto GetConvTestCases(miopenDataType_t datatype) +{ + using TestCase = miopen::unit_tests::ConvTestCase; + + return std::vector{ + // clang-format off + TestCase{{1, 1, 5, 10}, {16, 1, 5, 10}, {0, 0}, {2, 2}, {1, 1}, datatype}, + // clang-format on + }; +} + +const auto& GetTestParams() +{ + static const auto params = [] { + // gfx90A is not enabled because of WORKAROUND_ISSUE_1146 + Gpu supported_gpus = Gpu::gfx900 | Gpu::gfx906 | Gpu::gfx908; + auto p = miopen::unit_tests::UnitTestConvSolverParams(supported_gpus); + p.CheckXnackDisabled(); + return p; + }(); + return params; +} + +} // namespace + +using GPU_UnitTestConvSolverAsm5x10u2v2f1Fwd_FP32 = GPU_UnitTestConvSolverFwd_FP32; + +using CPU_UnitTestConvSolverAsm5x10u2v2f1DevApplicabilityFwd_NONE = + CPU_UnitTestConvSolverDevApplicabilityFwd_NONE; + +TEST_P(GPU_UnitTestConvSolverAsm5x10u2v2f1Fwd_FP32, ConvAsm5x10u2v2f1) +{ + this->RunTest(miopen::solver::conv::ConvAsm5x10u2v2f1{}); +}; + +TEST_P(CPU_UnitTestConvSolverAsm5x10u2v2f1DevApplicabilityFwd_NONE, ConvAsm5x10u2v2f1) +{ + this->RunTest(miopen::solver::conv::ConvAsm5x10u2v2f1{}); +}; + +// Smoke tests +INSTANTIATE_TEST_SUITE_P(Smoke, + GPU_UnitTestConvSolverAsm5x10u2v2f1Fwd_FP32, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(miopenConvolutionAlgoDirect), + testing::ValuesIn(GetConvTestCases(miopenFloat)))); + +// Device applicability test +INSTANTIATE_TEST_SUITE_P(Smoke, + CPU_UnitTestConvSolverAsm5x10u2v2f1DevApplicabilityFwd_NONE, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(GetConvTestCases(miopenFloat)[0]))); diff --git a/test/gtest/unit_conv_solver_ConvAsm7x7c3h224w224k64u2v2p3q3f1.cpp b/test/gtest/unit_conv_solver_ConvAsm7x7c3h224w224k64u2v2p3q3f1.cpp new file mode 100644 index 0000000000..daae84afc7 --- /dev/null +++ b/test/gtest/unit_conv_solver_ConvAsm7x7c3h224w224k64u2v2p3q3f1.cpp @@ -0,0 +1,85 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2025 Advanced Micro Devices, Inc. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in all + * copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + * SOFTWARE. + * + *******************************************************************************/ + +#include "unit_conv_solver.hpp" + +namespace { + +auto GetConvTestCases(miopenDataType_t datatype) +{ + using TestCase = miopen::unit_tests::ConvTestCase; + + return std::vector{ + // clang-format off + TestCase{{1, 3, 224, 224}, {64, 3, 7, 7}, {3, 3}, {2, 2}, {1, 1}, datatype}, + // clang-format on + }; +} + +const auto& GetTestParams() +{ + static const auto params = [] { + // gfx90A is not enabled because of WORKAROUND_ISSUE_1146 + Gpu supported_gpus = Gpu::gfx900 | Gpu::gfx906 | Gpu::gfx908; + auto p = miopen::unit_tests::UnitTestConvSolverParams(supported_gpus); + p.CheckXnackDisabled(); + return p; + }(); + return params; +} + +} // namespace + +using GPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1Fwd_FP32 = GPU_UnitTestConvSolverFwd_FP32; + +using CPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1DevApplicabilityFwd_NONE = + CPU_UnitTestConvSolverDevApplicabilityFwd_NONE; + +TEST_P(GPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1Fwd_FP32, + ConvAsm7x7c3h224w224k64u2v2p3q3f1) +{ + this->RunTest(miopen::solver::conv::ConvAsm7x7c3h224w224k64u2v2p3q3f1{}); +}; + +TEST_P(CPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1DevApplicabilityFwd_NONE, + ConvAsm7x7c3h224w224k64u2v2p3q3f1) +{ + this->RunTest(miopen::solver::conv::ConvAsm7x7c3h224w224k64u2v2p3q3f1{}); +}; + +// Smoke tests +INSTANTIATE_TEST_SUITE_P(Smoke, + GPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1Fwd_FP32, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(miopenConvolutionAlgoDirect), + testing::ValuesIn(GetConvTestCases(miopenFloat)))); + +// Device applicability test +INSTANTIATE_TEST_SUITE_P( + Smoke, + CPU_UnitTestConvSolverAsm7x7c3h224w224k64u2v2p3q3f1DevApplicabilityFwd_NONE, + testing::Combine(testing::Values(GetTestParams()), + testing::Values(GetConvTestCases(miopenFloat)[0])));