Merge branch 'lrn_fp16_to_fp32' into test_migraphx_test_runner

src/targets/gpu/lowering.cpp

@@ -164,7 +164,6 @@ struct miopen_apply
         add_extend_op("gather");
         add_extend_op("leaky_relu");
         add_extend_op("logsoftmax");
-        add_extend_op("lrn");
         add_extend_op("multinomial");
         add_extend_op("nonzero");
         add_extend_op("pad");
@@ -192,6 +191,7 @@ struct miopen_apply
         add_gemm_op<op::quant_dot>("quant_dot");
         add_if_op();
         add_loop_op();
+        add_lrn_op();
         add_neg_op();
         add_nms_op();
         add_quant_convolution_op();
@@ -562,6 +562,40 @@ struct miopen_apply
             return mod->replace_instruction(ins, gpu_out);
         });
     }
+
+    void add_lrn_op()
+    {
+        apply_map.emplace("lrn", [=](instruction_ref ins) {
+            auto s      = ins->get_shape();
+            auto in     = ins->inputs().front();
+            auto output = insert_allocation(ins, s);
+
+            auto type = s.type();
+            if(type == shape::half_type)
+            {
+                shape s32{shape::float_type, s.lens()};
+                auto cout32 = mod->insert_instruction(
+                    ins, make_op("hip::allocate", {{"shape", to_value(s32)}}));
+                auto cop32     = make_op("convert", {{"target_type", shape::float_type}});
+                auto convert32 = mod->insert_instruction(
+                    ins, make_op("gpu::convert", cop32.to_value()), in, cout32);
+                auto lout32 = mod->insert_instruction(
+                    ins, make_op("hip::allocate", {{"shape", to_value(s32)}}));
+                auto lrn32 = mod->insert_instruction(
+                    ins, make_op("gpu::lrn", ins->get_operator().to_value()), convert32, lout32);
+                auto cop16  = make_op("convert", {{"target_type", shape::half_type}});
+                auto lout16 = mod->insert_instruction(
+                    ins, make_op("gpu::convert", cop16.to_value()), lrn32, output);
+                return mod->replace_instruction(ins, lout16);
+            }
+            else
+            {
+                auto lrn16 = mod->insert_instruction(
+                    ins, make_op("gpu::lrn", ins->get_operator().to_value()), in, output);
+                return mod->replace_instruction(ins, lrn16);
+            }
+        });
+    }
 };
 
 void lowering::apply(module& m) const { miopen_apply{&m, this}.apply(); }

test/verify/test_mul_lrn.cpp

+
+#include "verify_program.hpp"
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/make_op.hpp>
+
+struct test_mul_lrn : verify_program<test_mul_lrn>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::half_type, {1, 3, 224, 224}};
+        auto x = mm->add_parameter("x", s);
+        std::vector<float> vec(s.elements(), 500.0f);
+        auto l  = mm->add_literal(migraphx::literal(s, vec));
+        auto xl = mm->add_instruction(migraphx::make_op("mul"), x, l);
+        auto y  = mm->add_instruction(migraphx::make_op("relu"), xl);
+        mm->add_instruction(
+            migraphx::make_op("lrn",
+                              {{"alpha", 0.0001}, {"beta", 0.75}, {"bias", 1.0}, {"size", 50}}),
+            y);
+        return p;
+    }
+};