Add check_context pass

src/include/migraph/check_context.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_CHECK_CONTEXT_HPP
+#define MIGRAPH_GUARD_RTGLIB_CHECK_CONTEXT_HPP
+
+#include <migraph/program.hpp>
+
+namespace migraph {
+
+template<class T>
+struct check_context
+{
+    struct op
+    {
+        std::string name() const { return "check_context"; }
+        shape compute_shape(std::vector<shape>) const { return {}; }
+        argument compute(context& ctx, shape, std::vector<argument>) const
+        {
+            T* x = any_cast<T>(&ctx);
+            if(x == nullptr)
+                MIGRAPH_THROW(std::string("Unexpected context type: ") + ctx.type_id().name());
+            return {};
+        }
+    };
+
+    std::string name() const { return "check_context"; }
+    void apply(program& p) const
+    {
+        p.insert_instruction(p.begin(), op{});
+    }
+};
+
+} // namespace migraph
+
+#endif

src/include/migraph/operators.hpp

@@ -491,20 +491,6 @@ struct outline
     argument compute(context&, shape, std::vector<argument>) const { return {s, nullptr}; }
 };
 
-template <class T>
-struct check_context
-{
-    std::string name() const { return "check_context"; }
-    shape compute_shape(std::vector<shape>) const { return {}; }
-    argument compute(context& ctx, shape, std::vector<argument>) const
-    {
-        T* x = any_cast<T>(&ctx);
-        if(x == nullptr)
-            MIGRAPH_THROW(std::string("Unexpected context type: ") + ctx.type_id().name());
-        return {};
-    }
-};
-
 } // namespace migraph
 
 #endif

src/targets/gpu/lowering.cpp

@@ -249,7 +249,6 @@ struct miopen_apply
 
     void apply()
     {
-        prog->insert_instruction(prog->begin(), check_context<context>{});
         for(auto it = prog->begin(); it != prog->end(); it++)
         {
             if(it->op.name() == "convolution")

src/targets/gpu/target.cpp

@@ -2,13 +2,14 @@
 #include <migraph/gpu/lowering.hpp>
 #include <migraph/gpu/write_literals.hpp>
 #include <migraph/gpu/context.hpp>
+#include <migraph/check_context.hpp>
 
 namespace migraph {
 namespace gpu {
 
 std::vector<pass> target::get_passes(migraph::context&) const
 {
-    return {lowering{}, write_literals{}};
+    return {lowering{}, write_literals{}, check_context<context>{}};
 }
 
 std::string target::name() const { return "miopen"; }