Merge branch 'develop' of https://github.com/ROCmSoftwarePlatform/AMDMIGraphX into scalar_parsing

src/include/migraphx/op/relu.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct relu : unary
+struct relu : unary<relu>
 {
-    std::string name() const { return "relu"; }
+    auto apply() const
+    {
+        return [](auto x) { return std::max(decltype(x){0}, x); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/sigmoid.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct sigmoid : unary
+struct sigmoid : unary<sigmoid>
 {
-    std::string name() const { return "sigmoid"; }
+    auto apply() const
+    {
+        return [](auto x) { return 1.f / (1.f + std::exp(-x)); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/sin.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct sin : unary
+struct sin : unary<sin>
 {
-    std::string name() const { return "sin"; }
+    auto apply() const
+    {
+        return [](auto x) { return std::sin(x); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/sinh.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct sinh : unary
+struct sinh : unary<sinh>
 {
-    std::string name() const { return "sinh"; }
+    auto apply() const
+    {
+        return [](auto x) { return std::sinh(x); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/sub.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct sub : binary
+struct sub : binary<sub>
 {
-    std::string name() const { return "sub"; }
+    auto apply() const
+    {
+        return [](auto x, auto y) { return x - y; };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/tan.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct tan : unary
+struct tan : unary<tan>
 {
-    std::string name() const { return "tan"; }
+    auto apply() const
+    {
+        return [](auto x) { return std::tan(x); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/tanh.hpp

@@ -17,9 +17,12 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct tanh : unary
+struct tanh : unary<tanh>
 {
-    std::string name() const { return "tanh"; }
+    auto apply() const
+    {
+        return [](auto x) { return std::tanh(x); };
+    }
 };
 
 } // namespace op

src/include/migraphx/op/unary.hpp

 #ifndef MIGRAPHX_GUARD_OPERATORS_UNARY_HPP
 #define MIGRAPHX_GUARD_OPERATORS_UNARY_HPP
 
-#include <array>
-#include <migraphx/operation.hpp>
-#include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/literal.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
-#include <cmath>
-#include <utility>
+#include <migraphx/op/name.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace op {
 
-struct unary
+template <class Derived>
+struct unary : op_name<Derived>
 {
     shape compute_shape(std::vector<shape> inputs) const
     {
         check_shapes{inputs}.has(1);
         return inputs.at(0);
     }
+    argument compute(const shape& output_shape, std::vector<argument> args) const
+    {
+        argument result{output_shape};
+        visit_all(result, args[0])([&](auto output, auto input) {
+            if(input.get_shape().standard())
+            {
+                std::transform(input.begin(),
+                               input.end(),
+                               output.begin(),
+                               static_cast<const Derived&>(*this).apply());
+            }
+            else
+            {
+                shape_for_each(output.get_shape(), [&](const auto& idx) {
+                    output(idx.begin(), idx.end()) =
+                        static_cast<const Derived&>(*this).apply()(input(idx.begin(), idx.end()));
+                });
+            }
+        });
+        return result;
+    }
 };
 
 } // namespace op

src/include/migraphx/program.hpp

@@ -108,6 +108,7 @@ struct program
     void debug_print() const;
     void debug_print(instruction_ref ins) const;
     void debug_print(const std::vector<instruction_ref>& inss) const;
+    void print_graph(std::ostream& os) const;
 
     void dry_run(parameter_map params) const;
 

src/include/migraphx/constant_propagate.hpp → src/include/migraphx/propagate_constant.hpp

-#ifndef MIGRAPHX_GUARD_RTGLIB_CONSTANT_PROPAGATE_HPP
-#define MIGRAPHX_GUARD_RTGLIB_CONSTANT_PROPAGATE_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_PROPAGATE_CONSTANT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_PROPAGATE_CONSTANT_HPP
 
 #include <string>
 #include <migraphx/config.hpp>
@@ -12,9 +12,9 @@ struct program;
 /**
  * Replace instructions which take all literals with a literal of the computation.
  */
-struct constant_propagate
+struct propagate_constant
 {
-    std::string name() const { return "constant_propagate"; }
+    std::string name() const { return "propagate_constant"; }
     void apply(program& p) const;
 };
 

src/include/migraphx/ranges.hpp

@@ -12,7 +12,7 @@ inline namespace MIGRAPHX_INLINE_NS {
 namespace detail {
 
 template <class String, class T>
-auto generic_find_impl(rank<2>, String&& s, const T& x) -> decltype(s.begin() + s.find(x), s.npos)
+auto generic_find_impl(rank<2>, String&& s, const T& x) -> decltype(s.npos, s.begin() + s.find(x))
 {
     auto index = s.find(x);
     if(index == s.npos)

src/program.cpp

@@ -56,7 +56,7 @@ static void print_instruction(std::ostream& os,
 }
 
 template <class F>
-static void print_program(std::ostream& os, const program& p, F annonate)
+static void print_program(const program& p, F print_func)
 {
     std::unordered_map<instruction_ref, std::string> names;
     int count = 0;
@@ -77,11 +77,7 @@ static void print_program(std::ostream& os, const program& p, F annonate)
             (void)arg;
         }
 
-        print_instruction(os, ins, names);
-
-        annonate(ins, names);
-
-        os << std::endl;
+        print_func(ins, names);
 
         count++;
     }
@@ -460,10 +456,12 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
     double calculate_overhead_time    = total_time - total_instruction_time;
     double calculate_overhead_percent = calculate_overhead_time * 100.0 / total_time;
 
-    print_program(os, *this, [&](auto ins, auto&&) {
+    print_program(*this, [&](auto ins, const auto& names) {
+        print_instruction(std::cout, ins, names);
         double avg     = common_average(ins_vec[ins]);
         double percent = std::ceil(100.0 * avg / total_instruction_time);
         os << ": " << avg << "ms, " << percent << "%";
+        os << std::endl;
     });
 
     os << std::endl;
@@ -501,7 +499,7 @@ void program::debug_print(instruction_ref ins) const
         return;
     }
     std::stringstream ss;
-    print_program(ss, *this, [&](auto x, auto&& names) {
+    print_program(*this, [&](auto x, const auto& names) {
         if(x == ins)
         {
             print_instruction(std::cout, x, names);
@@ -516,6 +514,32 @@ void program::debug_print(const std::vector<instruction_ref>& inss) const
     std::cout << std::endl;
 }
 
+static std::string enclose_name(const std::string& name)
+{
+    return '"' + replace_string(name, "\"", "\\\"") + '"';
+}
+
+void program::print_graph(std::ostream& os) const
+{
+    os << "digraph {" << std::endl;
+    os << "\trankdir=LR;" << std::endl;
+    print_program(*this, [&](auto ins, const auto& names) {
+        os << "\t" << enclose_name(names.at(ins))
+           << "[label=" << enclose_name(to_string(ins->get_operator())) << "];";
+        os << std::endl;
+        if(!ins->inputs().empty())
+        {
+            for(auto&& arg : ins->inputs())
+            {
+                os << "\t" << enclose_name(names.at(arg)) << " -> " << enclose_name(names.at(ins));
+                os << "[label=" << enclose_name(to_string(ins->get_shape())) << "];";
+                os << std::endl;
+            }
+        }
+    });
+    os << "}" << std::endl;
+}
+
 void program::dry_run(std::unordered_map<std::string, argument> params) const
 {
     auto& ctx = this->impl->ctx;
@@ -524,14 +548,21 @@ void program::dry_run(std::unordered_map<std::string, argument> params) const
 
 void program::annotate(std::ostream& os, std::function<void(instruction_ref)> a) const
 {
-    print_program(os, *this, [&](auto ins, auto&&) { a(ins); });
+    print_program(*this, [&](auto ins, const auto& names) {
+        print_instruction(os, ins, names);
+        a(ins);
+        os << std::endl;
+    });
 }
 
 bool operator==(const program& x, const program& y) { return to_string(x) == to_string(y); }
 
 std::ostream& operator<<(std::ostream& os, const program& p)
 {
-    print_program(os, p, [](auto&&...) {});
+    print_program(p, [&](auto ins, const auto& names) {
+        print_instruction(os, ins, names);
+        os << std::endl;
+    });
     return os;
 }
 

src/propagate_constant.cpp

+#include <migraphx/propagate_constant.hpp>
+#include <migraphx/program.hpp>
+#include <migraphx/matcher.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/functional.hpp>
+#include <unordered_set>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+bool skip_propogate(instruction_ref ins)
+{
+    if(ins->name() == "@literal")
+        return true;
+    auto&& s = ins->get_shape();
+    if(s.broadcasted() and not s.scalar())
+        return true;
+    if(s.scalar() and s.elements() != 1)
+        return true;
+    return false;
+}
+
+void propagate_constant::apply(program& p) const
+{
+    fix([&](auto self, auto ins) {
+        if(not skip_propogate(ins))
+        {
+            auto r = ins->eval();
+            if(not r.empty())
+            {
+                assert(r.get_shape() == ins->get_shape());
+                auto l = p.add_literal(r.get_shape(), r.data());
+                p.replace_instruction(ins, l);
+                return;
+            }
+        }
+        std::unordered_set<instruction_ref> children(ins->inputs().begin(), ins->inputs().end());
+        for(auto child : children)
+            self(child);
+    })(std::prev(p.end()));
+}
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/cpu/lowering.cpp

@@ -602,6 +602,7 @@ struct cpu_unary
                 std::transform(input.begin(), input.end(), output.begin(), op.fcn());
             });
         });
+
         return result;
     }
 };

src/targets/gpu/CMakeLists.txt

@@ -65,6 +65,7 @@ add_library(migraphx_gpu
     gather.cpp
     lrn.cpp
     schedule_model.cpp
+    adjust_allocation.cpp
 )
 set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
 rocm_clang_tidy_check(migraphx_gpu)

src/targets/gpu/abs.cpp

@@ -8,7 +8,7 @@ namespace gpu {
 shape miopen_abs::compute_shape(const std::vector<shape>& inputs) const
 {
     check_shapes{inputs, *this}.has(2).not_broadcasted();
-    return inputs.at(1);
+    return inputs.at(0);
 }
 
 argument miopen_abs::compute(context& ctx,

src/targets/gpu/adjust_allocation.cpp

+#include <migraphx/gpu/adjust_allocation.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/program.hpp>
+#include <migraphx/iterator_for.hpp>
+#include <algorithm>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+void adjust_allocation::apply(program& p) const
+{
+    for(auto ins : iterator_for(p))
+    {
+        // skip instruction with no input
+        if(ins->inputs().empty())
+            continue;
+
+        if(ins->name() == "load")
+            continue;
+
+        auto alias_ins = instruction::get_output_alias(ins, true);
+        if(alias_ins->name() == "hip::allocate")
+        {
+            // shape allocated is different from actual shape
+            // of the instruction, reallocate and replace the previous one
+            if(alias_ins->get_shape() != ins->get_shape())
+            {
+                auto alloc_ins = p.insert_instruction(ins, hip_allocate{ins->get_shape()});
+                p.replace_instruction(alias_ins, alloc_ins);
+            }
+        }
+    }
+}
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/include/migraphx/gpu/adjust_allocation.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_ADJUST_ALLOCATION_HPP
+#define MIGRAPHX_GUARD_RTGLIB_ADJUST_ALLOCATION_HPP
+
+#include <migraphx/program.hpp>
+#include <migraphx/config.hpp>
+#include <migraphx/gpu/context.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+
+namespace gpu {
+
+struct adjust_allocation
+{
+    std::string name() const { return "gpu::adjust_allocation"; }
+    void apply(program& p) const;
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/include/migraphx/gpu/oper.hpp

@@ -45,7 +45,7 @@ struct unary_device : oper<Derived>
     shape compute_shape(const std::vector<shape>& inputs) const
     {
         check_shapes{inputs, *this}.has(2);
-        return inputs.at(0);
+        return inputs.at(1);
     }
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const
@@ -63,7 +63,7 @@ struct binary_device : oper<Derived>
     shape compute_shape(const std::vector<shape>& inputs) const
     {
         check_shapes{inputs, *this}.has(3);
-        return inputs.at(0);
+        return inputs.at(2);
     }
 
     argument compute(context& ctx, const shape&, const std::vector<argument>& args) const

src/targets/gpu/tanh.cpp

@@ -8,7 +8,7 @@ namespace gpu {
 shape miopen_tanh::compute_shape(const std::vector<shape>& inputs) const
 {
     check_shapes{inputs, *this}.has(2).not_broadcasted();
-    return inputs.at(1);
+    return inputs.at(0);
 }
 
 argument miopen_tanh::compute(context& ctx,