[No Review Needed] test for CI docker build with ck commit hash change

requirements.txt

@@ -7,5 +7,5 @@ nlohmann/[email protected] -DJSON_MultipleHeaders=ON -DJSON_BuildTests=Off
 ROCm/[email protected]
 ROCm/[email protected]
 ROCm/frugally-deep@9683d557eb672ee2304f80f6682c51242d748a50
-ROCm/composable_kernel@421588131ef08dde36971bc1fa2016c005c9a0dd -DCMAKE_BUILD_TYPE=Release
+ROCm/composable_kernel@cedccd59c94cb0c74e7ec0d0f6c791aed081febc -DCMAKE_BUILD_TYPE=Release
 google/[email protected]

src/include/miopen/conv/solvers.hpp

@@ -5130,6 +5130,100 @@ struct ConvHipImplicitGemmF16F8F16WrwXdlops final
     bool CheckCKApplicability(const miopen::conv::ProblemDescription&) const;
 };
 
+struct PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops
+    : PerfConfigBaseCK<PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops>
+{
+    int index;
+    int split_k;
+    std::string kernel_id;
+    std::vector<std::string> valid_kernels;
+    PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops(int idx, std::string kernl_id)
+        : index(idx), kernel_id(kernl_id)
+    {
+    }
+    PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops()
+        : PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops(0, "")
+    {
+    }
+    PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops(bool)
+        : PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops(0, "")
+    {
+    }
+    MIOPEN_INTERNALS_EXPORT void HeuristicInit(const ExecutionContext&,
+                                               const miopen::conv::ProblemDescription&);
+    MIOPEN_INTERNALS_EXPORT bool SetNextValue(const miopen::conv::ProblemDescription&);
+    MIOPEN_INTERNALS_EXPORT bool IsValidValue() const;
+    bool IsValid(const ExecutionContext&, const miopen::conv::ProblemDescription& problem) const
+    {
+        return IsValid(problem);
+    }
+    MIOPEN_INTERNALS_EXPORT bool IsValid(const miopen::conv::ProblemDescription&) const;
+    MIOPEN_INTERNALS_EXPORT bool
+    operator==(const PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops& other) const;
+    MIOPEN_INTERNALS_EXPORT bool
+    IsModelApplicable(const ExecutionContext& ctx,
+                      const miopen::conv::ProblemDescription& problem) const;
+
+private:
+#if MIOPEN_ENABLE_AI_KERNEL_TUNING
+    std::vector<int> heuristic_indexes;
+    std::unordered_map<int, std::vector<std::string>> heuristic_kernels;
+    template <typename DataType>
+    bool RunParameterPredictionModel(const ExecutionContext& ctx,
+                                     const miopen::conv::ProblemDescription& problem);
+    void InitHeuristicKernelIDs(const std::string& type);
+    bool ModelApplyToken(int idx,
+                         std::string value,
+                         const std::string& arch,
+                         const miopen::conv::ProblemDescription& problem);
+#endif
+    template <typename DataType>
+    void Init(const miopen::conv::ProblemDescription&);
+    template <typename DataType>
+    bool CheckIsSupportCKArgs(const miopen::conv::ProblemDescription&) const;
+};
+
+struct ConvHipImplicitGemmGroupWrwNCHWXdlops final
+    : ConvTunableSolver<PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops>
+{
+    const std::string& SolverDbId() const override
+    {
+        return GetSolverDbId<ConvHipImplicitGemmGroupWrwNCHWXdlops>();
+    }
+
+    MIOPEN_INTERNALS_EXPORT PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops
+    GetDefaultPerformanceConfig(const ExecutionContext&,
+                                const miopen::conv::ProblemDescription&) const override;
+    MIOPEN_INTERNALS_EXPORT bool
+    IsValidPerformanceConfig(const ExecutionContext&,
+                             const miopen::conv::ProblemDescription&,
+                             const PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops&) const override;
+    MIOPEN_INTERNALS_EXPORT PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops
+    Search(const ExecutionContext&,
+           const miopen::conv::ProblemDescription&,
+           const AnyInvokeParams& invoke_ctx) const override;
+    MIOPEN_INTERNALS_EXPORT bool
+    IsApplicable(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
+    bool IsDynamic() const override { return true; }
+    MIOPEN_INTERNALS_EXPORT ConvSolution
+    GetSolution(const ExecutionContext&,
+                const miopen::conv::ProblemDescription&,
+                const PerformanceConfigHipImplicitGemmGroupWrwNCHWXdlops&) const override;
+    /// \ref igemm_get_wti_magic_number
+    float GetWti(const ExecutionContext&, const miopen::conv::ProblemDescription&) const override
+    {
+        return 0.02f;
+    };
+
+    MIOPEN_INTERNALS_EXPORT size_t GetWorkspaceSize(
+        const ExecutionContext&, const miopen::conv::ProblemDescription&) const override;
+    bool MayNeedWorkspace() const override { return true; }
+
+private:
+    template <typename DataType>
+    bool CheckCKApplicability(const miopen::conv::ProblemDescription&) const;
+};
+
 } // namespace conv
 } // namespace solver
 } // namespace miopen

src/mlo_dir_conv.cpp

@@ -156,6 +156,7 @@ static auto GetImplicitGemmWrWSolvers()
         miopen::solver::conv::ConvAsmImplicitGemmGTCDynamicWrwXdlops,
 #if MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL
         miopen::solver::conv::ConvHipImplicitGemmGroupWrwXdlops,
+        miopen::solver::conv::ConvHipImplicitGemmGroupWrwNCHWXdlops,
         miopen::solver::conv::ConvHipImplicitGemm3DGroupWrwXdlops,
         miopen::solver::conv::ConvHipImplicitGemmF16F8F16WrwXdlops,
 #endif // MIOPEN_BACKEND_HIP && MIOPEN_USE_COMPOSABLEKERNEL

src/solver.cpp

@@ -709,6 +709,11 @@ inline SolverRegistrar::SolverRegistrar(IdRegistryData& registry)
              multimarginloss::MultiMarginLossForward{}.SolverDbId());
 
     Register(registry, ++id, Primitive::Mha, mha::MhaCKFlashAttentionV2Forward{}.SolverDbId());
+
+    RegisterWithSolver(registry,
+                       ++id,
+                       conv::ConvHipImplicitGemmGroupWrwNCHWXdlops{},
+                       miopenConvolutionAlgoImplicitGEMM);
     // IMPORTANT: New solvers should be added to the end of the function, and don't leave a white
     // space between this comment and the newly registered solver(s)!
 }