From cbfa84bf11e5a3ce3778db2e387e3c075b510d1c Mon Sep 17 00:00:00 2001
From: Aleksandr Eremin <CAHEK7@yandex.ru>
Date: Tue, 12 Dec 2023 22:55:00 +0200
Subject: [PATCH 1/2] Properly guard CK usage by MIOPEN_USE_COMPOSABLEKERNEL
defines
---
src/solver/conv_hip_implicit_gemm_f16f8f16_fwd_xdlops.cpp | 5 +++--
src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp | 6 ++++--
2 files changed, 7 insertions(+), 4 deletions(-)
diff --git a/src/solver/conv_hip_implicit_gemm_f16f8f16_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_f16f8f16_fwd_xdlops.cpp
index dcbf668d80..8c8757e1d0 100644
--- a/src/solver/conv_hip_implicit_gemm_f16f8f16_fwd_xdlops.cpp
+++ b/src/solver/conv_hip_implicit_gemm_f16f8f16_fwd_xdlops.cpp
@@ -44,11 +44,12 @@ namespace miopen {
namespace solver {
namespace conv {
-using d_type = ck::half_t;
-using c_type = ck::f8_t;
using ProblemDescription = miopen::conv::ProblemDescription;
#if MIOPEN_USE_COMPOSABLEKERNEL
+using d_type = ck::half_t;
+using c_type = ck::f8_t;
+
template <typename DataType, typename ComputeType>
using DeviceOpF8Fwd = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD<
3,
diff --git a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
index e0a5f3921f..99f8c654e1 100644
--- a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
+++ b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
@@ -191,7 +191,6 @@ bool ConvHipImplicitGemmGroupFwdXdlops::CheckCKApplicability(
{
return IsCKApplicable<DeviceOpGFwdPtrs<DataType>, CKArgs>(problem);
}
-#endif
#if MIOPEN_ENABLE_AI_KERNEL_TUNING
static std::vector<std::string> GetKernelAsTokens(const std::string& kernel)
@@ -285,7 +284,8 @@ bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::RunParameterPredictionModel
}
return false;
}
-#endif
+#endif // MIOPEN_ENABLE_AI_KERNEL_TUNING
+#endif // MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::IsModelApplicable(
const ExecutionContext& ctx, const ProblemDescription& problem) const
@@ -339,6 +339,7 @@ void PerformanceConfigHipImplicitGemmGroupFwdXdlops::HeuristicInit(
bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::SetNextValue(const ProblemDescription& problem)
{
+#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
if(valid_kernels.empty())
{
switch(problem.GetInDataType())
@@ -362,6 +363,7 @@ bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::SetNextValue(const ProblemD
return true;
}
else
+#endif
return false;
}
From abd49ce6fb35ba4a679ff6d02cc4e96ecbbe10d9 Mon Sep 17 00:00:00 2001
From: Jun Liu <Liu.Jun@amd.com>
Date: Tue, 12 Dec 2023 17:52:18 -0800
Subject: [PATCH 2/2] Update
src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
Co-authored-by: Artem Tamazov <artem.tamazov@gmail.com>
---
src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
index 99f8c654e1..af8274db7b 100644
--- a/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
+++ b/src/solver/conv_hip_implicit_gemm_grouped_fwd_xdlops.cpp
@@ -339,7 +339,7 @@ void PerformanceConfigHipImplicitGemmGroupFwdXdlops::HeuristicInit(
bool PerformanceConfigHipImplicitGemmGroupFwdXdlops::SetNextValue(const ProblemDescription& problem)
{
-#if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
+#if MIOPEN_USE_COMPOSABLEKERNEL
if(valid_kernels.empty())
{
switch(problem.GetInDataType())