Merged from master ... still debugging ResNet18

Dockerfile

@@ -59,7 +59,7 @@ RUN ln -s $PREFIX /opt/rocm/hcc
 # Install dependencies
 ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
-RUN cget -p $PREFIX install -f /dev-requirements.txt
+RUN cget -p $PREFIX install -f /dev-requirements.txt -DMIOPEN_CACHE_DIR=""
 
 ENV LD_LIBRARY_PATH=$PREFIX/lib
 

requirements.txt

 google/protobuf -DCMAKE_POSITION_INDEPENDENT_CODE=On
 RadeonOpenCompute/rocm-cmake@3f43e2d493f24abbab4dc189a9ab12cc3ad33baf --build
-ROCmSoftwarePlatform/MIOpen,http://gitlab1.amd.com/pfultz/miopen/-/archive/57dd8372fa043f4bb500ac297f2c479e1c819e89/miopen-57dd8372fa043f4bb500ac297f2c479e1c819e89.tar.bz2 -DMIOPEN_CACHE_DIR=""
+ROCmSoftwarePlatform/MIOpen,http://gitlab1.amd.com/pfultz/miopen/-/archive/57dd8372fa043f4bb500ac297f2c479e1c819e89/miopen-57dd8372fa043f4bb500ac297f2c479e1c819e89.tar.bz2
 ROCmSoftwarePlatform/rocBLAS@v14.0.1
 blaze,https://bitbucket.org/blaze-lib/blaze/get/f0755dea0e03.tar.gz -X header -DHEADER_DIR=blaze

src/CMakeLists.txt

@@ -2,6 +2,8 @@
 add_library(migraph 
     auto_contiguous.cpp
     dead_code_elimination.cpp
+    eliminate_contiguous.cpp
+    env.cpp
     generate.cpp
     program.cpp
     shape.cpp

src/eliminate_contiguous.cpp

+#include <migraph/eliminate_contiguous.hpp>
+#include <migraph/program.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/ranges.hpp>
+#include <migraph/stringutils.hpp>
+
+namespace migraph {
+
+bool try_compute_shape(operation op, std::vector<instruction_ref> args)
+{
+    try
+    {
+        compute_shape(op, args);
+    }
+    catch(...)
+    {
+        return false;
+    }
+    return true;
+}
+
+void eliminate_contiguous::apply(program& p) const
+{
+    for(auto ins : iterator_for(p))
+    {
+        // Make a copy so we can modify it while we iterate
+        auto args = ins->arguments;
+        for(auto arg : ins->arguments)
+        {
+            // TODO: Pass in names for the operator in the constructor instead
+            // of using ends_with
+            if(ends_with(arg->op.name(), "contiguous"))
+            {
+                auto new_args = args;
+                auto prev     = arg->arguments.front();
+                replace(new_args, arg, prev);
+                if(try_compute_shape(ins->op, new_args))
+                {
+                    replace_argument(ins, arg, prev);
+                }
+            }
+        }
+    }
+}
+
+} // namespace migraph

src/env.cpp

+#include <migraph/env.hpp>
+#include <migraph/ranges.hpp>
+#include <cstdlib>
+
+namespace migraph {
+
+bool enabled(const char* name)
+{
+    auto e = env(name);
+    if(e.empty())
+        return false;
+    return contains({"1", "enable", "enabled", "yes", "true"}, e.front());
+}
+
+bool disabled(const char* name)
+{
+    auto e = env(name);
+    if(e.empty())
+        return false;
+    return contains({"0", "disable", "disabled", "no", "false"}, e.front());
+}
+
+std::vector<std::string> env(const char* name)
+{
+    auto p = std::getenv(name);
+    if(p == nullptr)
+        return {};
+    else
+        return {{p}};
+}
+
+} // namespace migraph

src/include/migraph/eliminate_contiguous.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_ELIMINATE_CONTIGUOUS_HPP
+#define MIGRAPH_GUARD_RTGLIB_ELIMINATE_CONTIGUOUS_HPP
+
+#include <string>
+#include <migraph/instruction_ref.hpp>
+
+namespace migraph {
+
+struct program;
+
+struct eliminate_contiguous
+{
+    std::string name() const { return "eliminate_contiguous"; }
+    void apply(program& p) const;
+};
+
+} // namespace migraph
+
+#endif

src/include/migraph/env.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_ENV_HPP
+#define MIGRAPH_GUARD_RTGLIB_ENV_HPP
+
+#include <vector>
+#include <string>
+
+namespace migraph {
+
+// Declare a cached environment variable
+#define MIGRAPH_DECLARE_ENV_VAR(x)                \
+    struct x                                      \
+    {                                             \
+        static const char* value() { return #x; } \
+    }; // NOLINT
+
+bool enabled(const char* name);
+bool disabled(const char* name);
+std::vector<std::string> env(const char* name);
+
+template <class T>
+bool enabled(T)
+{
+    static const bool result = enabled(T::value());
+    return result;
+}
+
+template <class T>
+bool disabled(T)
+{
+    static const bool result = disabled(T::value());
+    return result;
+}
+
+} // namespace migraph
+
+#endif

src/include/migraph/instruction.hpp

@@ -24,6 +24,7 @@ struct instruction
 
     instruction(literal l) : op(builtin::literal{}), result(l.get_shape()), lit(std::move(l)) {}
 
+    // internal
     void replace(operation o, shape r, std::vector<instruction_ref> args)
     {
         op = o;
@@ -46,12 +47,14 @@ struct instruction
 
     void recompute_shape() { replace(compute_shape(op, arguments)); }
 
+    // internal
     void replace(std::vector<instruction_ref> args)
     {
         clear_arguments();
         arguments = std::move(args);
     }
 
+    // internal
     void replace_argument(instruction_ref old, instruction_ref new_ins)
     {
         std::replace(arguments.begin(), arguments.end(), old, new_ins);

src/include/migraph/ranges.hpp

@@ -2,13 +2,62 @@
 #define MIGRAPH_GUARD_MIGRAPHLIB_RANGES_HPP
 
 #include <algorithm>
+#include <initializer_list>
 
 namespace migraph {
 
+template <int N>
+struct rank : rank<N - 1>
+{
+};
+
+template <>
+struct rank<0>
+{
+};
+
+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 index = s.find(x);
+    if(index == s.npos)
+        return s.end();
+    else
+        return s.begin() + index;
+}
+
 template <class C, class T>
-bool contains(C&& c, T&& x)
+auto generic_find_impl(rank<1>, C&& c, const T& x) -> decltype(c.find(x))
 {
-    return c.find(x) != c.end();
+    return c.find(x);
+}
+
+template <class C, class T>
+auto generic_find_impl(rank<0>, C&& c, const T& x)
+{
+    return std::find(c.begin(), c.end(), x);
+}
+
+} // namespace detail
+
+template <class C, class T>
+auto generic_find(C&& c, const T& x)
+{
+    return detail::generic_find_impl(rank<2>{}, c, x);
+}
+
+template <class C, class T>
+bool contains(const C& c, const T& x)
+{
+    return generic_find(c, x) != c.end();
+}
+
+template <class T, class U>
+bool contains(const std::initializer_list<T>& c, const U& x)
+{
+    return generic_find(c, x) != c.end();
 }
 
 template <class Range, class Iterator>
@@ -17,6 +66,12 @@ void copy(Range&& r, Iterator it)
     std::copy(r.begin(), r.end(), it);
 }
 
+template <class Range, class T>
+void replace(Range&& r, const T& old, const T& new_x)
+{
+    std::replace(r.begin(), r.end(), old, new_x);
+}
+
 template <class Iterator>
 struct iterator_range
 {

src/onnx/CMakeLists.txt

@@ -18,7 +18,7 @@ target_link_libraries(read_onnx migraph_onnx)
 
 add_executable(mnist mnist.cpp)
 rocm_clang_tidy_check(mnist)
-target_link_libraries(mnist migraph_cpu migraph_onnx)
+target_link_libraries(mnist migraph_cpu migraph_gpu migraph_onnx)
 
 add_executable(resnet18 resnet18.cpp)
 rocm_clang_tidy_check(resnet18)

src/onnx/mnist.cpp

@@ -6,7 +6,8 @@
 
 #include <migraph/onnx.hpp>
 
-#include <migraph/cpu/cpu_target.hpp>
+#include <migraph/gpu/target.hpp>
+#include <migraph/gpu/hip.hpp>
 #include <migraph/generate.hpp>
 
 auto reverse_int(unsigned int i)
@@ -120,15 +121,19 @@ int main(int argc, char const* argv[])
 
         std::string file = argv[1];
         auto prog        = migraph::parse_onnx(file);
-        prog.compile(migraph::cpu::cpu_target{});
+        std::cout << prog << std::endl << std::endl;
+        prog.compile(migraph::gpu::target{});
         auto s = migraph::shape{migraph::shape::float_type, {1, 1, 28, 28}};
         std::cout << s << std::endl;
         auto ptr = input.data();
+        migraph::program::parameter_map m;
+        m["output"] =
+            migraph::gpu::to_gpu(migraph::generate_argument(prog.get_parameter_shape("output")));
         for(int i = 0; i < 20; i++)
         {
             std::cout << "label: " << labels[i] << "  ---->  ";
-            auto input3 = migraph::argument{s, &ptr[784 * i]};
-            auto result = prog.eval({{"Input3", input3}});
+            m["0"]      = migraph::gpu::to_gpu(migraph::argument{s, &ptr[784 * i]});
+            auto result = migraph::gpu::from_gpu(prog.eval(m));
             std::vector<float> logits;
             result.visit([&](auto output) { logits.assign(output.begin(), output.end()); });
             std::vector<float> probs = softmax(logits);

src/onnx/onnx.cpp

@@ -286,7 +286,6 @@ struct onnx_parser
     {
         nodes = get_nodes(graph);
         std::unordered_map<std::string, onnx::TensorProto> initializer_data;
-        auto cnt = 0;
         for(auto&& f : graph.initializer())
         {
             initializer_data[f.name()] = f;

src/onnx/resnet18.cpp

@@ -18,14 +18,14 @@ auto read_cifar10_images(std::string full_path)
 
     const size_t nimages          = 10;
     const size_t nbytes_per_image = 3072;
-    std::vector<int8_t> raw_data(nimages * (nbytes_per_image + 1));
-    std::vector<int8_t> labels(nimages);
+    std::vector<uint8_t> raw_data(nimages * (nbytes_per_image + 1));
+    std::vector<uint8_t> labels(nimages);
     std::vector<float> data(nimages * nbytes_per_image);
     if(file.is_open())
     {
         file.read(reinterpret_cast<char*>(raw_data.data()),
-                  (nbytes_per_image + 1) * nimages * sizeof(int8_t));
-        int8_t* pimage = raw_data.data();
+                  (nbytes_per_image + 1) * nimages * sizeof(uint8_t));
+        uint8_t* pimage = raw_data.data();
         for(size_t i = 0; i < nimages; i++, pimage += nbytes_per_image)
         {
             labels[i] = *pimage++;
@@ -58,7 +58,7 @@ int main(int argc, char const* argv[])
     std::string file     = argv[1];
     std::string datafile = argv[2];
     auto prog            = migraph::parse_onnx(file);
-
+    std::cout << prog << std::endl;
     auto imageset = read_cifar10_images(datafile);
 
     // GPU target
@@ -81,6 +81,14 @@ int main(int argc, char const* argv[])
         for(auto x : logits)
             std::cout << x << "  ";
         std::cout << std::endl;
+        std::cout << std::endl;
+
+        for (int j = 0; j < 10; j++) {
+            std::cout << 255.0*input[i*3072+j] << "    ";
+        }
+        std::cout << std::endl;
+        std::cout << std::endl;
+        std::cout << std::endl;
     }
 
     // // // CPU target

src/program.cpp

 #include <migraph/program.hpp>
 #include <migraph/stringutils.hpp>
 #include <migraph/instruction.hpp>
+#include <migraph/env.hpp>
 #include <iostream>
 #include <sstream>
 #include <algorithm>
 
 namespace migraph {
 
+MIGRAPH_DECLARE_ENV_VAR(MIGRAPH_TRACE_COMPILE)
+
 struct program_impl
 {
     // A list is used to keep references to an instruction stable
@@ -183,9 +186,16 @@ void program::compile(const target& t)
 {
     assert(this->validate() == impl->instructions.end());
     this->impl->ctx = t.get_context();
+    if(enabled(MIGRAPH_TRACE_COMPILE{}))
+        std::cout << *this << std::endl << std::endl;
+    ;
     for(auto&& p : t.get_passes(this->impl->ctx))
     {
+        if(enabled(MIGRAPH_TRACE_COMPILE{}))
+            std::cout << "Pass: " << p.name() << std::endl;
         p.apply(*this);
+        if(enabled(MIGRAPH_TRACE_COMPILE{}))
+            std::cout << *this << std::endl << std::endl;
 #ifndef NDEBUG
         auto invalid = this->validate();
         if(invalid != impl->instructions.end())
@@ -230,7 +240,7 @@ argument program::eval(std::unordered_map<std::string, argument> params) const
                            ins.arguments.end(),
                            values.begin(),
                            [&](instruction_ref i) { return results.at(std::addressof(*i)); });
-            std::cout << "Compute: " << ins.op.name() << std::endl;
+            // std::cout << "Compute: " << ins.op.name() << std::endl;
             result = ins.op.compute(this->impl->ctx, ins.result, values);
         }
         results.emplace(std::addressof(ins), result);

src/targets/gpu/hip.cpp

@@ -36,7 +36,7 @@ std::vector<T> read_from_gpu(const void* x, std::size_t sz)
     std::vector<T> result(sz);
     auto status = hipMemcpy(result.data(), x, sz * sizeof(T), hipMemcpyDeviceToHost);
     if(status != hipSuccess)
-        MIGRAPH_THROW("Copy from gpu failed: " + hip_error(status));
+        MIGRAPH_THROW("Copy from gpu failed: " + hip_error(status)); // NOLINT
     return result;
 }
 

src/targets/gpu/target.cpp

@@ -6,6 +6,7 @@
 #include <migraph/auto_contiguous.hpp>
 #include <migraph/dead_code_elimination.hpp>
 #include <migraph/simplify_reshapes.hpp>
+#include <migraph/eliminate_contiguous.hpp>
 
 namespace migraph {
 namespace gpu {
@@ -16,11 +17,16 @@ std::vector<pass> target::get_passes(migraph::context& gctx) const
     // clang-format off
     return
     {
+        //dead_code_elimination{},
         auto_contiguous{},
-        simplify_reshapes{},
+        //simplify_reshapes{},
+        //dead_code_elimination{},
         lowering{ctx},
+        //eliminate_contiguous{},
+        //dead_code_elimination{},
+        //write_literals{},
         write_literals{},
-        check_context<context>{},
+        //check_context<context>{},
         dead_code_elimination{}
     };
     // clang-format on

test/auto_contiguous_test.cpp

@@ -13,23 +13,6 @@ struct contiguous_target
     migraph::context get_context() const { return {}; }
 };
 
-migraph::literal get_2x2()
-{
-    return migraph::literal{{migraph::shape::float_type, {2, 2}}, {1, 2, 3, 4}};
-}
-
-migraph::literal get_2x2_transposed()
-{
-    return migraph::literal{{migraph::shape::float_type, {2, 2}, {1, 2}}, {1, 2, 3, 4}};
-}
-
-migraph::literal get_2() { return migraph::literal{{migraph::shape::float_type, {2}}, {1, 2}}; }
-
-migraph::literal get_2_broadcasted()
-{
-    return migraph::literal{{migraph::shape::float_type, {2, 1}, {1, 0}}, {1, 2}};
-}
-
 void literal_broadcast()
 {
     migraph::program p;

test/eliminate_contiguous_test.cpp

+#include <migraph/eliminate_contiguous.hpp>
+#include <migraph/dead_code_elimination.hpp>
+#include <migraph/operators.hpp>
+#include <basic_ops.hpp>
+#include <test.hpp>
+
+struct eliminate_contiguous_target
+{
+    std::string name() const { return "eliminate_contiguous"; }
+    std::vector<migraph::pass> get_passes(migraph::context&) const
+    {
+        return {migraph::eliminate_contiguous{}, migraph::dead_code_elimination{}};
+    }
+    migraph::context get_context() const { return {}; }
+};
+
+void standard_op()
+{
+    migraph::program p;
+    auto l = p.add_literal(get_2x2());
+    auto t = p.add_instruction(migraph::transpose{{1, 0}}, l);
+    auto c = p.add_instruction(migraph::contiguous{}, t);
+    p.add_instruction(pass_standard_op{}, c);
+    auto count = std::distance(p.begin(), p.end());
+    p.compile(eliminate_contiguous_target{});
+    EXPECT(std::distance(p.begin(), p.end()) == count);
+}
+
+void non_standard_op()
+{
+    migraph::program p;
+    auto l = p.add_literal(get_2x2());
+    auto t = p.add_instruction(migraph::transpose{{1, 0}}, l);
+    auto c = p.add_instruction(migraph::contiguous{}, t);
+    p.add_instruction(pass_op{}, c);
+    auto count = std::distance(p.begin(), p.end());
+    p.compile(eliminate_contiguous_target{});
+    EXPECT(std::distance(p.begin(), p.end()) == (count - 1));
+}
+
+int main()
+{
+    standard_op();
+    non_standard_op();
+}

test/gpu/miopen.cpp

@@ -11,6 +11,9 @@
 
 #include <miopen/miopen.h>
 
+#include <future>
+#include <thread>
+
 #include "test.hpp"
 #include "verify.hpp"
 
@@ -19,6 +22,17 @@
 #pragma clang diagnostic ignored "-Wglobal-constructors"
 #endif
 
+// An improved async, that doesn't block
+template <class Function>
+std::future<typename std::result_of<Function()>::type> detach_async(Function&& f)
+{
+    using result_type = typename std::result_of<Function()>::type;
+    std::packaged_task<result_type()> task(std::forward<Function>(f));
+    auto fut = task.get_future();
+    std::thread(std::move(task)).detach();
+    return std::move(fut);
+}
+
 struct auto_print
 {
     static std::array<std::function<void()>, 2> handlers;
@@ -85,9 +99,9 @@ void verify_program()
         for(auto&& handle : auto_print::handlers)
             handle();
     });
-    auto cpu_arg = run_cpu<V>();
-    auto gpu_arg = run_gpu<V>();
-    visit_all(cpu_arg, gpu_arg)([](auto cpu, auto gpu) {
+    auto cpu_arg_f = detach_async([] { return run_cpu<V>(); });
+    auto gpu_arg   = run_gpu<V>();
+    visit_all(cpu_arg_f.get(), gpu_arg)([](auto cpu, auto gpu) {
         if(not test::verify_range(cpu, gpu))
         {
             std::cout << "FAILED: " << migraph::get_type_name<V>() << std::endl;

test/include/basic_ops.hpp

@@ -81,6 +81,30 @@ struct pass_op
     }
 };
 
+struct pass_standard_op
+{
+    std::string name() const { return "pass"; }
+    migraph::argument
+    compute(migraph::context&, migraph::shape, std::vector<migraph::argument> args) const
+    {
+        if(args.empty())
+            return {};
+        return args.front();
+    }
+
+    migraph::shape compute_shape(std::vector<migraph::shape> inputs) const
+    {
+        for(auto&& input : inputs)
+        {
+            if(not input.standard())
+                throw std::runtime_error("Not standard shape");
+        }
+        if(inputs.empty())
+            return {};
+        return inputs.front();
+    }
+};
+
 struct nop
 {
     std::string name() const { return "nop"; }
@@ -92,3 +116,23 @@ struct nop
 
     migraph::shape compute_shape(std::vector<migraph::shape>) const { return {}; }
 };
+
+inline migraph::literal get_2x2()
+{
+    return migraph::literal{{migraph::shape::float_type, {2, 2}}, {1, 2, 3, 4}};
+}
+
+inline migraph::literal get_2x2_transposed()
+{
+    return migraph::literal{{migraph::shape::float_type, {2, 2}, {1, 2}}, {1, 2, 3, 4}};
+}
+
+inline migraph::literal get_2()
+{
+    return migraph::literal{{migraph::shape::float_type, {2}}, {1, 2}};
+}
+
+inline migraph::literal get_2_broadcasted()
+{
+    return migraph::literal{{migraph::shape::float_type, {2, 1}, {1, 0}}, {1, 2}};
+}