Add tracer

src/include/migraph/generate.hpp

@@ -10,7 +10,7 @@ namespace migraph {
 template <class T>
 struct xorshf96_generator
 {
-    unsigned long max = 31;
+    long max = 16;
     unsigned long x   = 123456789;
     unsigned long y   = 362436069;
     unsigned long z   = 521288629;
@@ -26,7 +26,7 @@ struct xorshf96_generator
         y               = z;
         z               = t ^ x ^ y;
 
-        return z % max;
+        return z % max - (max + 1);
     }
 };
 

src/include/migraph/program.hpp

@@ -8,6 +8,7 @@
 #include <migraph/builtin.hpp>
 #include <migraph/instruction_ref.hpp>
 #include <migraph/target.hpp>
+#include <migraph/tracer.hpp>
 #include <algorithm>
 #include <iostream>
 
@@ -88,7 +89,7 @@ struct program
 
     instruction_ref validate() const;
 
-    void compile(const target& t);
+    void compile(const target& t, tracer trace = tracer{});
 
     void perf_report(std::ostream& os, std::size_t n, parameter_map params) const;
 

src/include/migraph/tensor_view.hpp

@@ -94,13 +94,13 @@ struct tensor_view
     // TODO: Add iterators so it can handle nonstandard tensors
     T* begin()
     {
-        assert(this->m_shape.standard());
+        assert(this->m_shape.standard() or this->empty());
         return m_data;
     }
 
     T* end()
     {
-        assert(this->m_shape.standard());
+        assert(this->m_shape.standard() or this->empty());
         if(this->empty())
             return m_data;
         else
@@ -109,13 +109,13 @@ struct tensor_view
 
     const T* begin() const
     {
-        assert(this->m_shape.standard());
+        assert(this->m_shape.standard() or this->empty());
         return m_data;
     }
 
     const T* end() const
     {
-        assert(this->m_shape.standard());
+        assert(this->m_shape.standard() or this->empty());
         if(this->empty())
             return m_data;
         else

src/include/migraph/tracer.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_TRACER_HPP
+#define MIGRAPH_GUARD_RTGLIB_TRACER_HPP
+
+#include <ostream>
+
+namespace migraph {
+
+struct swallow
+{
+    template<class... Ts>
+    swallow(Ts&&...)
+    {}
+};
+
+struct tracer
+{
+    tracer()
+    {}
+
+    tracer(std::ostream& s) : os(&s)
+    {}
+
+    bool enabled() const
+    {
+        return os != nullptr;
+    }
+
+    template<class... Ts>
+    void operator()(const Ts&... xs) const
+    {
+        if(os != nullptr) 
+        {
+            swallow{*os << xs...};
+            *os << std::endl;
+        }
+    }
+
+private:
+    std::ostream * os = nullptr;
+};
+
+} // namespace migraph
+
+#endif

src/program.cpp

@@ -237,23 +237,21 @@ instruction_ref program::validate() const
                         [&](const instruction& i) { return !i.valid(impl->instructions.begin()); });
 }
 
-void program::compile(const target& t)
+void program::compile(const target& t, tracer trace)
 {
     assert(this->validate() == impl->instructions.end());
     this->impl->ctx = t.get_context();
-    if(enabled(MIGRAPH_TRACE_COMPILE{}))
-        std::cout << *this << std::endl << std::endl;
-    ;
+    if(not trace.enabled() and enabled(MIGRAPH_TRACE_COMPILE{}))
+        trace = tracer{std::cout};
+    trace(*this);
+    trace();
     for(auto&& p : t.get_passes(this->impl->ctx))
     {
-        if(enabled(MIGRAPH_TRACE_COMPILE{}))
-            std::cout << "Pass: " << p.name() << std::endl;
+        trace("Pass: ", p.name());
         p.apply(*this);
-        if(enabled(MIGRAPH_TRACE_COMPILE{}))
-            std::cout << *this << std::endl;
+        trace(*this);
 #ifndef NDEBUG
-        if(enabled(MIGRAPH_TRACE_COMPILE{}))
-            std::cout << "Validate ..." << std::endl;
+        trace("Validate ...");
         auto invalid = this->validate();
         if(invalid != impl->instructions.end())
         {
@@ -261,8 +259,7 @@ void program::compile(const target& t)
             MIGRAPH_THROW(p.name() + " pass produces invalid program at instruction " +
                           std::to_string(index) + ": " + invalid->op.name());
         }
-        if(enabled(MIGRAPH_TRACE_COMPILE{}))
-            std::cout << std::endl;
+        trace();
 #endif
     }
     auto invalid = this->validate();

src/targets/cpu/cpu_lowering.cpp

@@ -212,6 +212,7 @@ struct cpu_contiguous
     shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
     argument compute(context&, const shape& output_shape, std::vector<argument> args) const
     {
+        assert(output_shape.standard());
         argument result{output_shape};
         visit_all(result, args[0])([&](auto output, auto input) {
             shape_for_each(output.get_shape(), [&](const auto& idx) {

src/targets/gpu/lowering.cpp

@@ -250,6 +250,8 @@ struct miopen_contiguous
     }
     argument compute(context&, shape output_shape, const std::vector<argument>& args) const
     {
+        assert(output_shape == args[1].get_shape());
+        assert(output_shape.standard());
         hip_contiguous(std::move(output_shape), args.at(0), args.at(1));
         return args.at(1);
     }

test/gpu/miopen.cpp

@@ -24,13 +24,17 @@
 
 // An improved async, that doesn't block
 template <class Function>
-std::future<typename std::result_of<Function()>::type> detach_async(Function&& f)
+std::future<typename std::result_of<Function()>::type> detach_async(Function&& f, bool parallel=true)
 {
+    if(parallel){   
     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);
+    } else {
+    return std::async(std::launch::deferred, std::move(f));
+    }
 }
 
 struct auto_print
@@ -50,13 +54,26 @@ struct auto_print
 };
 std::array<std::function<void()>, 2> auto_print::handlers = {};
 
+void compile_check(migraph::program& p, migraph::target t)
+{
+    auto name = t.name();
+    auto s = p.get_shape();
+    std::stringstream ss;
+    p.compile(std::move(t), migraph::tracer{ss});
+    if(p.get_shape() != s) 
+    {
+        std::cout << ss.str() << std::endl;
+        throw std::runtime_error("Compiling program with " + name + " alters its shape");
+    }
+}
+
 template <class V>
 migraph::argument run_cpu()
 {
     V v;
     auto p = v.create_program();
     auto_print pp{p, 0};
-    p.compile(migraph::cpu::cpu_target{});
+    compile_check(p, migraph::cpu::cpu_target{});
     migraph::program::parameter_map m;
     for(auto&& x : p.get_parameter_shapes())
     {
@@ -71,7 +88,7 @@ migraph::argument run_gpu()
     V v;
     auto p = v.create_program();
     auto_print pp{p, 1};
-    p.compile(migraph::gpu::target{});
+    compile_check(p, migraph::gpu::target{});
 
     migraph::program::parameter_map m;
     for(auto&& x : p.get_parameter_shapes())
@@ -255,6 +272,7 @@ struct test_contiguous
         migraph::shape s{migraph::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
         auto x = p.add_parameter("x", s);
         p.add_instruction(migraph::contiguous{}, x);
+        EXPECT(p.get_shape().standard());
         return p;
     }
 };