Merge branch 'master' into memory_coloring

src/targets/gpu/eliminate_allocation.cpp

@@ -20,7 +20,7 @@ void eliminate_allocation::apply(program& p) const
             continue;
         allocs.emplace_back(ins, n);
         std::size_t size = ins->get_shape().bytes();
-        n += size + (size % 4);
+        n += size + (size % 32);
     }
     auto mem = p.add_parameter("memory", shape{shape::int8_type, {n}});
     for(auto&& pp : allocs)

src/targets/gpu/fuse_ops.cpp

@@ -12,7 +12,7 @@ struct hip_add_relu
     std::string name() const { return "hip::add_relu"; }
     shape compute_shape(const std::vector<shape>& inputs) const
     {
-        check_shapes{inputs}.has(3).standard();
+        check_shapes{inputs, *this}.has(3);
         return inputs.front();
     }
     argument compute(context&, const shape&, const std::vector<argument>& args) const

src/targets/gpu/include/migraph/gpu/device/add.hpp

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_ADD_HPP
+
+#include <migraph/argument.hpp>
+
+namespace migraph {
+namespace gpu {
+namespace device {
+
+void add(const argument& result, const argument& arg1, const argument& arg2);
+
+} // namespace device
+} // namespace gpu
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/device/add_relu.hpp

@@ -8,7 +8,7 @@ namespace migraph {
 namespace gpu {
 namespace device {
 
-void add_relu(argument result, argument arg1, argument arg2);
+void add_relu(const argument& result, const argument& arg1, const argument& arg2);
 
 } // namespace device
 } // namespace gpu

src/targets/gpu/lowering.cpp

@@ -9,6 +9,7 @@
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
 #include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
 #include <migraph/iterator_for.hpp>
 #include <migraph/gpu/rocblas.hpp>
 #include <migraph/gpu/context.hpp>
@@ -170,6 +171,23 @@ struct miopen_pooling
     }
 };
 
+struct hip_add
+{
+    std::string name() const { return "gpu::add"; }
+    shape compute_shape(const std::vector<shape>& inputs) const
+    {
+        // check_shapes{inputs, *this}.has(3).standard();
+        check_shapes{inputs, *this}.has(3);
+        return inputs.at(0);
+    }
+
+    argument compute(context&, const shape&, const std::vector<argument>& args) const
+    {
+        device::add(args[2], args[0], args[1]);
+        return args[2];
+    }
+};
+
 struct miopen_add
 {
     std::string name() const { return "gpu::add"; }
@@ -392,7 +410,7 @@ struct miopen_apply
     {
         auto output = insert_allocation(ins, ins->result);
         return prog->replace_instruction(
-            ins, miopen_add{}, ins->arguments.at(0), ins->arguments.at(1), output);
+            ins, hip_add{}, ins->arguments.at(0), ins->arguments.at(1), output);
     }
 
     instruction_ref apply_gemm(instruction_ref ins)

src/targets/gpu/target.cpp

@@ -11,6 +11,7 @@
 #include <migraph/dead_code_elimination.hpp>
 #include <migraph/simplify_reshapes.hpp>
 #include <migraph/eliminate_contiguous.hpp>
+#include <migraph/fwd_conv_batchnorm_rewrite.hpp>
 
 namespace migraph {
 namespace gpu {
@@ -21,6 +22,8 @@ std::vector<pass> target::get_passes(migraph::context& gctx) const
     // clang-format off
     return
     {
+        dead_code_elimination{},
+        fwd_conv_batchnorm_rewrite{},
         dead_code_elimination{},
         auto_contiguous{},
         simplify_reshapes{},

test/gpu/miopen.cpp

@@ -8,7 +8,7 @@
 #include <migraph/gpu/hip.hpp>
 #include <migraph/manage_ptr.hpp>
 #include <migraph/type_name.hpp>
-#include <migraph/verify.hpp>
+#include <migraph/verify_args.hpp>
 
 #include <miopen/miopen.h>
 
@@ -77,6 +77,12 @@ struct auto_print
 };
 std::array<std::function<void()>, 2> auto_print::handlers = {};
 
+template <class T>
+auto get_hash(const T& x)
+{
+    return std::hash<T>{}(x);
+}
+
 void compile_check(migraph::program& p, const migraph::target& t)
 {
     auto name = t.name();
@@ -100,7 +106,7 @@ migraph::argument run_cpu()
     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));
     }
     return p.eval(m);
 }
@@ -112,52 +118,15 @@ migraph::argument run_gpu()
     auto p = v.create_program();
     auto_print pp{p, 1};
     compile_check(p, 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)));
     }
 
     return migraph::gpu::from_gpu(p.eval(m));
 }
 
-void verify_args(const std::string& name,
-                 const migraph::argument& cpu_arg,
-                 const migraph::argument& gpu_arg)
-{
-    visit_all(cpu_arg, gpu_arg)([&](auto cpu, auto gpu) {
-        if(not migraph::verify_range(cpu, gpu))
-        {
-            // TODO: Check for nans
-            std::cout << "FAILED: " << name << std::endl;
-            // std::cout << cpu << std::endl;
-            // std::cout << gpu << std::endl;
-            if(migraph::range_zero(cpu))
-                std::cout << "Cpu data is all zeros" << std::endl;
-            if(migraph::range_zero(gpu))
-                std::cout << "Gpu data is all zeros" << std::endl;
-
-            auto idx = migraph::mismatch_idx(cpu, gpu, migraph::float_equal);
-            if(idx < migraph::range_distance(cpu))
-            {
-                std::cout << "Mismatch at " << idx << ": " << cpu[idx] << " != " << gpu[idx]
-                          << std::endl;
-            }
-
-            auto cpu_nan_idx = find_idx(cpu, migraph::not_finite);
-            if(cpu_nan_idx >= 0)
-                std::cout << "Non finite number found in cpu at " << cpu_nan_idx << ": "
-                          << cpu[cpu_nan_idx] << std::endl;
-
-            auto gpu_nan_idx = find_idx(gpu, migraph::not_finite);
-            if(gpu_nan_idx >= 0)
-                std::cout << "Non finite number found in gpu at " << gpu_nan_idx << ": "
-                          << gpu[gpu_nan_idx] << std::endl;
-        }
-    });
-}
-
 template <class V>
 void verify_program()
 {
@@ -210,6 +179,62 @@ struct test_add_broadcast
     }
 };
 
+struct test_add_broadcast2
+{
+    migraph::program create_program() const
+    {
+        migraph::program p;
+        migraph::shape s{migraph::shape::float_type, {3}};
+        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 3, 4}});
+        auto y  = p.add_parameter("y", {migraph::shape::float_type, {3}});
+        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
+        p.add_instruction(migraph::add{}, x, by);
+        return p;
+    }
+};
+
+struct test_add_broadcast3
+{
+    migraph::program create_program() const
+    {
+        migraph::program p;
+        migraph::shape s{migraph::shape::float_type, {3}};
+        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 4, 5}});
+        auto y  = p.add_parameter("y", {migraph::shape::float_type, {4}});
+        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
+        p.add_instruction(migraph::add{}, x, by);
+        return p;
+    }
+};
+
+struct test_add_broadcast4
+{
+    migraph::program create_program() const
+    {
+        migraph::program p;
+        migraph::shape s{migraph::shape::float_type, {3}};
+        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 3, 5}});
+        auto y  = p.add_parameter("y", {migraph::shape::float_type, {3}});
+        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
+        p.add_instruction(migraph::add{}, x, by);
+        return p;
+    }
+};
+
+struct test_add_broadcast5
+{
+    migraph::program create_program() const
+    {
+        migraph::program p;
+        migraph::shape s{migraph::shape::float_type, {3}};
+        auto x  = p.add_parameter("x", {migraph::shape::float_type, {2, 4, 8}});
+        auto y  = p.add_parameter("y", {migraph::shape::float_type, {4}});
+        auto by = p.add_instruction(migraph::broadcast{1}, x, y);
+        p.add_instruction(migraph::add{}, x, by);
+        return p;
+    }
+};
+
 struct test_conv_relu
 {
     migraph::program create_program() const
@@ -418,6 +443,10 @@ int main()
 {
     verify_program<test_add>();
     verify_program<test_add_broadcast>();
+    verify_program<test_add_broadcast2>();
+    verify_program<test_add_broadcast3>();
+    verify_program<test_add_broadcast4>();
+    verify_program<test_add_broadcast5>();
     verify_program<test_conv_relu>();
     verify_program<test_add_relu>();
     verify_program<test_conv_pooling>();