From cbfa84bf11e5a3ce3778db2e387e3c075b510d1c Mon Sep 17 00:00:00 2001 From: Aleksandr Eremin 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 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, CKArgs>(problem); } -#endif #if MIOPEN_ENABLE_AI_KERNEL_TUNING static std::vector 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 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 --- 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())