Merge branch 'master' into memory_coloring

src/include/migraph/generate.hpp

@@ -12,7 +12,7 @@ constexpr T normalize(unsigned long z)
 {
     if(z == 0)
         return 0;
-    const auto max     = 32768;
+    const auto max     = 2048;
     const double range = max / 2; // NOLINT
     double result      = (z % max) / range;
     result -= 1;

src/include/migraph/verify.hpp

@@ -162,10 +162,12 @@ double rms_range(R1&& r1, R2&& r2)
 }
 
 template <class R1, class R2>
-bool verify_range(R1&& r1, R2&& r2, double tolerance = 80)
+bool verify_range(R1&& r1, R2&& r2, double tolerance = 80, double* out_error = nullptr)
 {
     double threshold = std::numeric_limits<range_value<R1>>::epsilon() * tolerance;
     auto error       = rms_range(r1, r2);
+    if(out_error != nullptr)
+        *out_error = error;
     return error <= threshold;
 }
 } // namespace migraph

src/include/migraph/verify_args.hpp

@@ -12,10 +12,12 @@ inline void verify_args(const std::string& name,
                         double tolerance = 80)
 {
     visit_all(cpu_arg, gpu_arg)([&](auto cpu, auto gpu) {
-        if(not verify_range(cpu, gpu, tolerance))
+        double error;
+        if(not verify_range(cpu, gpu, tolerance, &error))
         {
             // TODO: Check for nans
             std::cout << "FAILED: " << name << std::endl;
+            std::cout << "error: " << error << std::endl;
             if(cpu.size() < 32)
                 std::cout << "cpu:" << cpu << std::endl;
             if(gpu.size() < 32)

src/onnx/verify_onnx.cpp

@@ -6,46 +6,109 @@
 #include <migraph/gpu/hip.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/verify_args.hpp>
+#include <migraph/instruction.hpp>
 
-migraph::argument run_cpu(const std::string& file)
+template <class T>
+auto get_hash(const T& x)
 {
-    auto p = migraph::parse_onnx(file);
+    return std::hash<T>{}(x);
+}
+
+template <class F>
+migraph::argument run_cpu(F f)
+{
+    auto p = f();
     p.compile(migraph::cpu::cpu_target{});
     migraph::program::parameter_map m;
     for(auto&& x : p.get_parameter_shapes())
     {
-        m[x.first] = migraph::generate_argument(x.second);
+        m[x.first] = migraph::generate_argument(x.second, get_hash(x.first));
     }
     auto out = p.eval(m);
     std::cout << p << std::endl;
     return out;
 }
 
-migraph::argument run_gpu(const std::string& file)
+template <class F>
+migraph::argument run_gpu(F f)
 {
-    auto p = migraph::parse_onnx(file);
+    auto p = f();
     p.compile(migraph::gpu::target{});
 
     migraph::program::parameter_map m;
     for(auto&& x : p.get_parameter_shapes())
     {
-        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second));
+        m[x.first] = migraph::gpu::to_gpu(migraph::generate_argument(x.second, get_hash(x.first)));
     }
     auto out = migraph::gpu::from_gpu(p.eval(m));
     std::cout << p << std::endl;
     return migraph::gpu::from_gpu(out);
 }
 
+template <class F>
+void verify_program(const std::string& name, F f, double tolerance = 100)
+{
+    auto x = run_cpu(f);
+    auto y = run_gpu(f);
+    migraph::verify_args(name, x, y, tolerance);
+}
+
+void verify_instructions(const migraph::program& prog, double tolerance = 80)
+{
+    for(auto&& ins : prog)
+    {
+        if(ins.op.name().front() == '@')
+            continue;
+        if(ins.op.name() == "broadcast")
+            continue;
+        if(ins.op.name() == "transpose")
+            continue;
+        if(ins.op.name() == "reshape")
+            continue;
+        auto create_program = [&] {
+            migraph::program p;
+            std::vector<migraph::instruction_ref> inputs;
+            for(auto&& arg : ins.arguments)
+            {
+                if(arg->op.name() == "@literal")
+                    inputs.push_back(p.add_literal(arg->lit));
+                else
+                    inputs.push_back(
+                        p.add_parameter(std::to_string(inputs.size()), arg->get_shape()));
+            }
+            p.add_instruction(ins.op, inputs);
+            return p;
+        };
+        try
+        {
+            std::cout << "Verify: " << ins.op.name() << std::endl;
+            std::cout << create_program() << std::endl;
+            verify_program(ins.op.name(), create_program, tolerance);
+        }
+        catch(...)
+        {
+            std::cout << "Instruction " << ins.op.name() << " threw an exception." << std::endl;
+            throw;
+        }
+    }
+}
+
 int main(int argc, char const* argv[])
 {
-    if(argc > 1)
+    std::vector<std::string> args(argv + 1, argv + argc);
+    if(not args.empty())
     {
-        std::string file = argv[1];
+        std::string file = args.front();
         auto p           = migraph::parse_onnx(file);
         std::cout << p << std::endl;
 
-        auto x = run_cpu(file);
-        auto y = run_gpu(file);
-        migraph::verify_args(file, x, y, 100);
+        if(std::any_of(args.begin(), args.end(), [](const auto& s) { return s == "-i"; }))
+        {
+            verify_instructions(p);
+        }
+        else
+        {
+            verify_program(file, [&] { return migraph::parse_onnx(file); });
+        }
     }
 }

src/program.cpp

@@ -179,6 +179,7 @@ instruction_ref program::add_outline(const shape& s)
 
 instruction_ref program::add_parameter(std::string name, shape s)
 {
+    assert(get_parameter_shape(name) == shape{});
     impl->instructions.push_front({builtin::param{std::move(name)}, std::move(s), {}});
     return impl->instructions.begin();
 }

src/targets/gpu/lowering.cpp

@@ -222,7 +222,7 @@ struct miopen_add
 struct miopen_gemm
 {
     gemm op;
-    std::string name() const { return "gpu::convolution"; }
+    std::string name() const { return "gpu::gemm"; }
     shape compute_shape(const std::vector<shape>& inputs) const
     {
         check_shapes{inputs, *this}.has(3);
@@ -357,7 +357,7 @@ struct miopen_apply
 
     instruction_ref insert_allocation(instruction_ref ins, const shape& s, std::string tag = "")
     {
-        if(ins == --prog->end())
+        if(ins == --prog->end() and not tag.empty())
         {
             return prog->add_parameter("output", s);
         }

test/gpu/miopen.cpp

@@ -439,6 +439,41 @@ struct test_conv_bn_relu_pooling
     }
 };
 
+struct test_conv_bn_relu_pooling2
+{
+    static migraph::instruction_ref
+    add_bn(migraph::program& p, migraph::instruction_ref x, std::size_t channels)
+    {
+        migraph::shape vars{migraph::shape::float_type, {channels}};
+        auto scale    = p.add_literal(migraph::abs(migraph::generate_literal(vars, 1 + channels)));
+        auto bias     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 2 + channels)));
+        auto mean     = p.add_literal(migraph::abs(migraph::generate_literal(vars, 3 + channels)));
+        auto variance = p.add_literal(migraph::abs(migraph::generate_literal(vars, 4 + channels)));
+        return p.add_instruction(migraph::batch_norm_inference{}, x, scale, bias, mean, variance);
+    }
+    migraph::program create_program() const
+    {
+        migraph::program p;
+
+        migraph::shape xs1{migraph::shape::float_type, {1, 512, 7, 7}};
+        migraph::shape xs2{migraph::shape::float_type, {1, 1024, 14, 14}};
+        migraph::shape ws1{migraph::shape::float_type, {2048, 512, 1, 1}};
+        migraph::shape ws2{migraph::shape::float_type, {2048, 1024, 1, 1}};
+        auto x1    = p.add_parameter("x1", xs1);
+        auto w1    = p.add_parameter("w1", ws1);
+        auto conv1 = p.add_instruction(migraph::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
+        auto bn1   = add_bn(p, conv1, 2048);
+        auto x2    = p.add_parameter("x2", xs2);
+        auto w2    = p.add_parameter("w2", ws2);
+        auto conv2 = p.add_instruction(migraph::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
+        auto bn2   = add_bn(p, conv2, 2048);
+        auto add   = p.add_instruction(migraph::add{}, bn1, bn2);
+        auto relu  = p.add_instruction(migraph::activation{"relu"}, add);
+        p.add_instruction(migraph::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
+        return p;
+    }
+};
+
 int main()
 {
     verify_program<test_add>();
@@ -460,4 +495,5 @@ int main()
     verify_program<test_batchnorm_inference>();
     verify_program<test_batchnorm_inference_2>();
     verify_program<test_conv_bn_relu_pooling>();
+    verify_program<test_conv_bn_relu_pooling2>();
 }