Enable pass

src/targets/gpu/CMakeLists.txt

@@ -93,6 +93,7 @@ add_library(migraphx_gpu
     compile_miopen.cpp
     compiler.cpp
     device_name.cpp
+    fuse_ck.cpp
     fuse_mlir.cpp
     fuse_ops.cpp
     gather.cpp

src/targets/gpu/fuse_ck.cpp

@@ -6,9 +6,6 @@
 #include <migraphx/env.hpp>
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
-
-MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_CK);
-
 struct module;
 
 namespace gpu {
@@ -122,11 +119,8 @@ struct find_ck_gemm
 
 void fuse_ck::apply(module_pass_manager& mpm) const
 {
-    if(enabled(MIGRAPHX_ENABLE_CK{}))
-    {
     match::find_matches(mpm, find_ck_gemm_pointwise{});
     match::find_matches(mpm, find_ck_gemm{});
-    }
 }
 
 } // namespace gpu

src/targets/gpu/target.cpp

@@ -57,6 +57,7 @@
 #include <migraphx/gpu/concat_gpu_opt.hpp>
 #include <migraphx/gpu/context.hpp>
 #include <migraphx/gpu/device_name.hpp>
+#include <migraphx/gpu/fuse_ck.hpp>
 #include <migraphx/gpu/fuse_mlir.hpp>
 #include <migraphx/gpu/fuse_ops.hpp>
 #include <migraphx/gpu/prefuse_ops.hpp>
@@ -74,6 +75,7 @@ namespace gpu {
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_SCHEDULE_PASS)
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_DISABLE_REDUCE_FUSION)
 MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_NHWC)
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_ENABLE_CK)
 struct id_pass
 {
     std::string name() const { return "id"; }
@@ -133,6 +135,8 @@ std::vector<pass> target::get_passes(migraphx::context& gctx, const compile_opti
         dead_code_elimination{},
         enable_pass(not enabled(MIGRAPHX_DISABLE_REDUCE_FUSION{}), fuse_reduce{}),
         dead_code_elimination{},
+        enable_pass(enabled(MIGRAPHX_ENABLE_CK{}), fuse_ck{}),
+        dead_code_elimination{},
         enable_pass(mlir_enabled(), fuse_mlir{&ctx}),
         dead_code_elimination{},
         lowering{&ctx, options.offload_copy},