Add support for builtin models in the driver (#404)

* Add experimental support for c++ output

* Format

* Fix syntax errors

* Add resnet50 model

* Formatting

* Add inceptionv3 model

* Formatting

* Add alexnet

* Formatting

* Fix name of pooling mode

* Formatting

* Fix tidy issues

* Ignore driver directory

* Show accetable values

CMakeLists.txt

@@ -55,8 +55,20 @@ elseif(MIGRAPHX_ENABLE_GPU)
 endif()
 rocm_enable_clang_tidy(
     CHECKS
-        *
-        -android-cloexec-fopen
+        boost-*
+        bugprone-*
+        cert-*
+        clang-analyzer-*
+        clang-diagnostic-*
+        cppcoreguidelines-*
+        google-*
+        hicpp-multiway-paths-covered
+        hicpp-signed-bitwise
+        llvm-namespace-comment
+        misc-*
+        modernize-*
+        performance-*
+        readability-*
         -clang-analyzer-alpha.core.CastToStruct
         -clang-analyzer-optin.performance.Padding
         -clang-diagnostic-deprecated-declarations
@@ -72,21 +84,9 @@ rocm_enable_clang_tidy(
         -cppcoreguidelines-pro-type-union-access
         -cppcoreguidelines-pro-type-vararg
         -cppcoreguidelines-special-member-functions
-        -fuchsia-*
-        -google-readability-braces-around-statements
-        -google-readability-todo
+        -google-readability-*
         -google-runtime-int
         -google-runtime-references
-        -hicpp-braces-around-statements
-        -hicpp-explicit-conversions
-        -hicpp-member-init
-        -hicpp-no-array-decay
-        -hicpp-special-member-functions
-        -hicpp-uppercase-literal-suffix
-        -hicpp-use-override
-        # This check is broken
-        -llvm-header-guard
-        -llvm-include-order
         -misc-macro-parentheses
         -modernize-concat-nested-namespaces
         -modernize-pass-by-value

codecov.yml

 ignore:
   - "test/"
+  - "src/driver"
  

src/driver/CMakeLists.txt

 
-add_executable(driver main.cpp verify.cpp perf.cpp)
+add_executable(driver 
+    main.cpp
+    verify.cpp
+    perf.cpp
+    resnet50.cpp
+    inceptionv3.cpp
+    alexnet.cpp
+)
 rocm_clang_tidy_check(driver)
 target_link_libraries(driver migraphx_cpu migraphx_onnx migraphx_tf)
 if(MIGRAPHX_ENABLE_GPU)

src/driver/alexnet.cpp

+#include <migraphx/operators.hpp>
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include "models.hpp"
+
+namespace migraphx {
+namespace driver {
+inline namespace MIGRAPHX_INLINE_NS {
+
+migraphx::program alexnet(unsigned batch) // NOLINT(readability-function-size)
+{
+    migraphx::program p;
+    auto m0 =
+        p.add_parameter("0", migraphx::shape{migraphx::shape::float_type, {batch, 3, 224, 224}});
+    auto mx0 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {1000}}, 0));
+    auto mx1 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {1000, 4096}}, 1));
+    auto mx2 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {4096}}, 2));
+    auto mx3 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {4096, 4096}}, 3));
+    auto mx4 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {4096}}, 4));
+    auto mx5 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {4096, 9216}}, 5));
+    auto mx6 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {256}}, 6));
+    auto mx7 = p.add_literal(migraphx::generate_literal(
+        migraphx::shape{migraphx::shape::float_type, {256, 256, 3, 3}}, 7));
+    auto mx8 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {256}}, 8));
+    auto mx9  = p.add_literal(migraphx::generate_literal(
+        migraphx::shape{migraphx::shape::float_type, {256, 384, 3, 3}}, 9));
+    auto mx10 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {384}}, 10));
+    auto mx11 = p.add_literal(migraphx::generate_literal(
+        migraphx::shape{migraphx::shape::float_type, {384, 192, 3, 3}}, 11));
+    auto mx12 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {192}}, 12));
+    auto mx13 = p.add_literal(migraphx::generate_literal(
+        migraphx::shape{migraphx::shape::float_type, {192, 64, 5, 5}}, 13));
+    auto mx14 = p.add_literal(
+        migraphx::generate_literal(migraphx::shape{migraphx::shape::float_type, {64}}, 14));
+    auto mx15 = p.add_literal(migraphx::generate_literal(
+        migraphx::shape{migraphx::shape::float_type, {64, 3, 11, 11}}, 15));
+    migraphx::op::convolution convolution16;
+    convolution16.padding  = {2, 2};
+    convolution16.stride   = {4, 4};
+    convolution16.dilation = {1, 1};
+    convolution16.group    = 1;
+    auto mx16              = p.add_instruction(convolution16, m0, mx15);
+    migraphx::op::broadcast broadcast17;
+    broadcast17.axis           = 1;
+    broadcast17.broadcast_lens = {batch, 64, 55, 55};
+    auto mx17                  = p.add_instruction(broadcast17, mx14);
+    migraphx::op::add add18;
+    auto mx18 = p.add_instruction(add18, mx16, mx17);
+    migraphx::op::relu relu19;
+    auto mx19 = p.add_instruction(relu19, mx18);
+    migraphx::op::pooling pooling20;
+    pooling20.mode    = "max";
+    pooling20.padding = {0, 0};
+    pooling20.stride  = {2, 2};
+    pooling20.lengths = {3, 3};
+    auto mx20         = p.add_instruction(pooling20, mx19);
+    migraphx::op::convolution convolution21;
+    convolution21.padding  = {2, 2};
+    convolution21.stride   = {1, 1};
+    convolution21.dilation = {1, 1};
+    convolution21.group    = 1;
+    auto mx21              = p.add_instruction(convolution21, mx20, mx13);
+    migraphx::op::broadcast broadcast22;
+    broadcast22.axis           = 1;
+    broadcast22.broadcast_lens = {batch, 192, 27, 27};
+    auto mx22                  = p.add_instruction(broadcast22, mx12);
+    migraphx::op::add add23;
+    auto mx23 = p.add_instruction(add23, mx21, mx22);
+    migraphx::op::relu relu24;
+    auto mx24 = p.add_instruction(relu24, mx23);
+    migraphx::op::pooling pooling25;
+    pooling25.mode    = "max";
+    pooling25.padding = {0, 0};
+    pooling25.stride  = {2, 2};
+    pooling25.lengths = {3, 3};
+    auto mx25         = p.add_instruction(pooling25, mx24);
+    migraphx::op::convolution convolution26;
+    convolution26.padding  = {1, 1};
+    convolution26.stride   = {1, 1};
+    convolution26.dilation = {1, 1};
+    convolution26.group    = 1;
+    auto mx26              = p.add_instruction(convolution26, mx25, mx11);
+    migraphx::op::broadcast broadcast27;
+    broadcast27.axis           = 1;
+    broadcast27.broadcast_lens = {batch, 384, 13, 13};
+    auto mx27                  = p.add_instruction(broadcast27, mx10);
+    migraphx::op::add add28;
+    auto mx28 = p.add_instruction(add28, mx26, mx27);
+    migraphx::op::relu relu29;
+    auto mx29 = p.add_instruction(relu29, mx28);
+    migraphx::op::convolution convolution30;
+    convolution30.padding  = {1, 1};
+    convolution30.stride   = {1, 1};
+    convolution30.dilation = {1, 1};
+    convolution30.group    = 1;
+    auto mx30              = p.add_instruction(convolution30, mx29, mx9);
+    migraphx::op::broadcast broadcast31;
+    broadcast31.axis           = 1;
+    broadcast31.broadcast_lens = {batch, 256, 13, 13};
+    auto mx31                  = p.add_instruction(broadcast31, mx8);
+    migraphx::op::add add32;
+    auto mx32 = p.add_instruction(add32, mx30, mx31);
+    migraphx::op::relu relu33;
+    auto mx33 = p.add_instruction(relu33, mx32);
+    migraphx::op::convolution convolution34;
+    convolution34.padding  = {1, 1};
+    convolution34.stride   = {1, 1};
+    convolution34.dilation = {1, 1};
+    convolution34.group    = 1;
+    auto mx34              = p.add_instruction(convolution34, mx33, mx7);
+    migraphx::op::broadcast broadcast35;
+    broadcast35.axis           = 1;
+    broadcast35.broadcast_lens = {batch, 256, 13, 13};
+    auto mx35                  = p.add_instruction(broadcast35, mx6);
+    migraphx::op::add add36;
+    auto mx36 = p.add_instruction(add36, mx34, mx35);
+    migraphx::op::relu relu37;
+    auto mx37 = p.add_instruction(relu37, mx36);
+    migraphx::op::pooling pooling38;
+    pooling38.mode    = "max";
+    pooling38.padding = {0, 0};
+    pooling38.stride  = {2, 2};
+    pooling38.lengths = {3, 3};
+    auto mx38         = p.add_instruction(pooling38, mx37);
+    migraphx::op::flatten flatten39;
+    flatten39.axis = 1;
+    auto mx39      = p.add_instruction(flatten39, mx38);
+    migraphx::op::identity identity40;
+    auto mx40 = p.add_instruction(identity40, mx39);
+    migraphx::op::transpose transpose41;
+    transpose41.dims = {1, 0};
+    auto mx41        = p.add_instruction(transpose41, mx5);
+    migraphx::op::multibroadcast multibroadcast42;
+    multibroadcast42.output_lens = {batch, 4096};
+    auto mx42                    = p.add_instruction(multibroadcast42, mx4);
+    migraphx::op::dot dot43;
+    dot43.alpha = 1;
+    dot43.beta  = 1;
+    auto mx43   = p.add_instruction(dot43, mx40, mx41, mx42);
+    migraphx::op::relu relu44;
+    auto mx44 = p.add_instruction(relu44, mx43);
+    migraphx::op::identity identity45;
+    auto mx45 = p.add_instruction(identity45, mx44);
+    migraphx::op::transpose transpose46;
+    transpose46.dims = {1, 0};
+    auto mx46        = p.add_instruction(transpose46, mx3);
+    migraphx::op::multibroadcast multibroadcast47;
+    multibroadcast47.output_lens = {batch, 4096};
+    auto mx47                    = p.add_instruction(multibroadcast47, mx2);
+    migraphx::op::dot dot48;
+    dot48.alpha = 1;
+    dot48.beta  = 1;
+    auto mx48   = p.add_instruction(dot48, mx45, mx46, mx47);
+    migraphx::op::relu relu49;
+    auto mx49 = p.add_instruction(relu49, mx48);
+    migraphx::op::transpose transpose50;
+    transpose50.dims = {1, 0};
+    auto mx50        = p.add_instruction(transpose50, mx1);
+    migraphx::op::multibroadcast multibroadcast51;
+    multibroadcast51.output_lens = {batch, 1000};
+    auto mx51                    = p.add_instruction(multibroadcast51, mx0);
+    migraphx::op::dot dot52;
+    dot52.alpha = 1;
+    dot52.beta  = 1;
+    p.add_instruction(dot52, mx49, mx50, mx51);
+    return p;
+}
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace driver
+} // namespace migraphx

src/driver/argument_parser.hpp

@@ -221,6 +221,11 @@ struct argument_parser
         return [=](auto&, auto& arg) { arg.metavar = metavar; };
     }
 
+    MIGRAPHX_DRIVER_STATIC auto type(const std::string& type)
+    {
+        return [=](auto&, auto& arg) { arg.type = type; };
+    }
+
     template <class T>
     MIGRAPHX_DRIVER_STATIC auto set_value(T value)
     {

src/driver/main.cpp

@@ -2,6 +2,7 @@
 #include "command.hpp"
 #include "verify.hpp"
 #include "perf.hpp"
+#include "models.hpp"
 
 #include <migraphx/tf.hpp>
 #include <migraphx/onnx.hpp>
@@ -26,17 +27,21 @@ inline namespace MIGRAPHX_INLINE_NS {
 
 struct loader
 {
+    std::string model;
     std::string file;
     std::string file_type;
-    bool is_nhwc  = true;
-    unsigned trim = 0;
-    bool optimize = false;
+    unsigned batch = 1;
+    bool is_nhwc   = true;
+    unsigned trim  = 0;
+    bool optimize  = false;
 
     void parse(argument_parser& ap)
     {
         ap(file, {}, ap.metavar("<input file>"));
+        ap(model, {"--model"}, ap.help("Load model"), ap.type("resnet50|inceptionv3|alexnet"));
         ap(file_type, {"--onnx"}, ap.help("Load as onnx"), ap.set_value("onnx"));
         ap(file_type, {"--tf"}, ap.help("Load as tensorflow"), ap.set_value("tf"));
+        ap(batch, {"--batch"}, ap.help("Set batch size for model"));
         ap(is_nhwc, {"--nhwc"}, ap.help("Treat tensorflow format as nhwc"), ap.set_value(true));
         ap(is_nhwc, {"--nchw"}, ap.help("Treat tensorflow format as nchw"), ap.set_value(false));
         ap(trim, {"--trim", "-t"}, ap.help("Trim instructions from the end"));
@@ -46,18 +51,32 @@ struct loader
     program load()
     {
         program p;
-        if(file_type.empty())
+        if(model.empty())
         {
-            if(ends_with(file, ".onnx"))
-                file_type = "onnx";
-            else if(ends_with(file, ".pb"))
-                file_type = "tf";
+            if(file_type.empty())
+            {
+                if(ends_with(file, ".onnx"))
+                    file_type = "onnx";
+                else if(ends_with(file, ".pb"))
+                    file_type = "tf";
+            }
+            std::cout << "Reading: " << file << std::endl;
+            if(file_type == "onnx")
+                p = parse_onnx(file);
+            else if(file_type == "tf")
+                p = parse_tf(file, is_nhwc);
+        }
+        else
+        {
+            if(model == "resnet50")
+                p = resnet50(batch);
+            else if(model == "inceptionv3")
+                p = inceptionv3(batch);
+            else if(model == "alexnet")
+                p = alexnet(batch);
+            else
+                MIGRAPHX_THROW("Unknown model: " + model);
         }
-        std::cout << "Reading: " << file << std::endl;
-        if(file_type == "onnx")
-            p = parse_onnx(file);
-        else if(file_type == "tf")
-            p = parse_tf(file, is_nhwc);
         if(trim > 0)
         {
             auto last = std::prev(p.end(), trim);
@@ -138,6 +157,7 @@ struct compiler
 struct read : command<read>
 {
     loader l;
+    bool cpp      = false;
     bool graphviz = false;
     bool brief    = false;
     std::string output;
@@ -149,6 +169,7 @@ struct read : command<read>
            ap.help("Print out a graphviz representation."),
            ap.set_value(true));
         ap(brief, {"--brief"}, ap.help("Make the output brief."), ap.set_value(true));
+        ap(cpp, {"--cpp"}, ap.help("Print out the program as cpp program."), ap.set_value(true));
         ap(output, {"--output", "-o"}, ap.help("Output to file."));
     }
 
@@ -164,7 +185,9 @@ struct read : command<read>
             os = &fs;
         }
 
-        if(graphviz)
+        if(cpp)
+            p.print_cpp(*os);
+        else if(graphviz)
             p.print_graph(*os, brief);
         else
             *os << p << std::endl;

src/driver/models.hpp

+
+#include <migraphx/program.hpp>
+
+namespace migraphx {
+namespace driver {
+inline namespace MIGRAPHX_INLINE_NS {
+
+migraphx::program resnet50(unsigned batch);
+migraphx::program inceptionv3(unsigned batch);
+migraphx::program alexnet(unsigned batch);
+
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace driver
+} // namespace migraphx

src/include/migraphx/program.hpp

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

src/include/migraphx/shape.hpp

@@ -181,11 +181,12 @@ struct shape
 #undef MIGRAPHX_SHAPE_GENERATE_VISITOR_ALL
     }
 
+    std::string type_string() const;
+
     private:
     std::shared_ptr<const shape_impl> impl;
 
     std::size_t element_space() const;
-    std::string type_string() const;
 };
 
 } // namespace MIGRAPHX_INLINE_NS

src/include/migraphx/stringutils.hpp

@@ -5,6 +5,7 @@
 #include <numeric>
 #include <string>
 #include <sstream>
+#include <vector>
 #include <migraphx/config.hpp>
 
 namespace migraphx {
@@ -43,6 +44,18 @@ inline std::string join_strings(Strings strings, const std::string& delim)
     });
 }
 
+inline std::vector<std::string> split_string(const std::string& s, char delim)
+{
+    std::vector<std::string> elems;
+    std::stringstream ss(s + ' ');
+    std::string item;
+    while(std::getline(ss, item, delim))
+    {
+        elems.push_back(item);
+    }
+    return elems;
+}
+
 template <class F>
 std::string trim(const std::string& s, F f)
 {

src/program.cpp

@@ -624,6 +624,104 @@ void program::print_graph(std::ostream& os, bool brief) const
     os << "}" << std::endl;
 }
 
+static std::string cpp_var_name(const std::string& name)
+{
+    return "m" + replace_string(name, "@", "x");
+}
+
+static std::string cpp_op_var(const std::string& name, instruction_ref ins)
+{
+    return replace_string(name, "@", ins->name());
+}
+
+static void print_op_attributes(std::ostream& os, const std::string& name, const operation& op)
+{
+    std::string x = to_string(op);
+    if(contains(x, "["))
+    {
+        auto start                 = x.find('[');
+        auto end                   = x.find(']');
+        std::string attribute_text = x.substr(start + 1, end - start - 1);
+        std::vector<std::string> attributes;
+        for(auto&& attribute : split_string(attribute_text, ','))
+        {
+            if(contains(attribute, '='))
+                attributes.push_back(attribute);
+            else
+                attributes.back() += "," + attribute;
+        }
+        for(auto&& attribute : attributes)
+        {
+            auto p     = split_string(attribute, '=');
+            auto key   = p.front();
+            auto value = p.back();
+            if(contains({"bn_mode", "padding_mode"}, key))
+                continue;
+            if(key == "mode")
+                value = enclose_name(trim(value));
+            os << name << "." << key << " = " << value << ";" << std::endl;
+        }
+    }
+}
+
+static void print_cpp_shape(std::ostream& os, const migraphx::shape& s)
+{
+    os << "migraphx::shape{migraphx::shape::" << s.type_string();
+    os << ", {" << to_string_range(s.lens()) << "}";
+    if(not s.standard())
+        os << ", {" << to_string_range(s.strides()) << "}";
+    os << "}";
+}
+
+void program::print_cpp(std::ostream& os) const
+{
+    os << "migraphx::program p;" << std::endl;
+    // cppcheck-suppress variableScope
+    unsigned long seed = 0;
+    print_program(*this, [&](auto ins, const auto& names) {
+        auto op = cpp_op_var(names.at(ins), ins);
+        if(ins->name().front() != '@')
+        {
+            os << "migraphx::op::" << ins->name() << " " << op << ";" << std::endl;
+            print_op_attributes(os, op, ins->get_operator());
+        }
+        os << "auto " << cpp_var_name(names.at(ins)) << " = ";
+        if(ins->name() == "@literal")
+        {
+            os << "p.add_literal(";
+            bool use_abs = false;
+            ins->get_literal().visit([&](auto v) {
+                use_abs = std::none_of(v.begin(), v.end(), [](auto x) { return x < 0; });
+            });
+            if(use_abs)
+                os << "migraphx::abs(";
+            os << "migraphx::generate_literal(";
+            print_cpp_shape(os, ins->get_shape());
+            os << ", " << seed << ")";
+            if(use_abs)
+                os << ")";
+            os << ");" << std::endl;
+            seed++;
+        }
+        else if(ins->name() == "@param")
+        {
+            std::string name = any_cast<builtin::param>(ins->get_operator()).parameter;
+            os << "p.add_parameter(" << enclose_name(name) << ",";
+            print_cpp_shape(os, ins->get_shape());
+            os << ");" << std::endl;
+        }
+        else
+        {
+            os << "p.add_instruction(" << op;
+            for(auto input : ins->inputs())
+            {
+                os << ", " << cpp_var_name(names.at(input));
+            }
+            os << ");" << std::endl;
+        }
+    });
+}
+
 void program::dry_run(std::unordered_map<std::string, argument> params) const
 {
     auto& ctx = this->impl->ctx;

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

@@ -99,8 +99,10 @@ inline pooling_descriptor make_pooling(const migraphx::op::pooling& op)
     miopenPoolingMode_t mode;
     if(op.mode == "max")
         mode = miopenPoolingMax;
-    else
+    else if(op.mode == "average")
         mode = miopenPoolingAverage;
+    else
+        MIGRAPHX_THROW("Unknown mode for pooling: " + op.mode);
     auto p = make_obj<pooling_descriptor>(&miopenCreatePoolingDescriptor);
     miopenSet2dPoolingDescriptor(p.get(),
                                  mode,