Write literals to the gpu

src/targets/miopen/include/migraph/miopen/hip.hpp

@@ -26,6 +26,20 @@ struct hip_allocate
     }
 };
 
+struct hip_write
+{
+    std::string name() const { return "hip::write"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return inputs.front();
+    }
+    argument compute(context&, shape, std::vector<argument> args) const
+    {
+        return to_gpu(args.front());
+    }
+};
+
 } // namespace miopen
 
 } // namespace migraph

src/targets/miopen/miopen_target.cpp

@@ -6,6 +6,7 @@
 #include <migraph/miopen/miopen.hpp>
 #include <migraph/miopen/hip.hpp>
 #include <migraph/dfor.hpp>
+#include <migraph/iterator_for.hpp>
 
 namespace migraph {
 namespace miopen {
@@ -306,7 +307,25 @@ struct miopen_pass
     void apply(program& p) const { miopen_apply{&p}.apply(); }
 };
 
-std::vector<pass> miopen_target::get_passes(context&) const { return {miopen_pass{}}; }
+struct miopen_write_literals
+{
+    std::string name() const { return "miopen::write_literals"; }
+
+    void apply(program& p) const 
+    { 
+        for(auto ins:iterator_for(p))
+        {
+            if(ins->op.name() == "@literal")
+            {
+                literal l = ins->lit;
+                auto pre = p.add_literal(l);
+                p.replace_instruction(ins, hip_write{}, pre);
+            }
+        }
+    }
+};
+
+std::vector<pass> miopen_target::get_passes(context&) const { return {miopen_pass{}, miopen_write_literals{}}; }
 
 std::string miopen_target::name() const { return "miopen"; }