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

Jenkinsfile

@@ -101,7 +101,8 @@ rocmtest tidy: rocmnode('rocmtest') { cmake_build ->
     stage('Clang Debug') {
         // TODO: Enanle integer
         def sanitizers = "undefined"
-        cmake_build("hcc", "-DCMAKE_BUILD_TYPE=debug -DCMAKE_CXX_FLAGS_DEBUG='-g -fno-omit-frame-pointer -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}'")
+        def debug_flags = "-g -fno-omit-frame-pointer -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
+        cmake_build("hcc", "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'")
     }
     stage('Clang Release') {
         cmake_build("hcc", "-DCMAKE_BUILD_TYPE=release")
@@ -120,7 +121,7 @@ rocmtest tidy: rocmnode('rocmtest') { cmake_build ->
         // TODO: Add bounds-strict
         def sanitizers = "undefined,address"
         def debug_flags = "-g -fprofile-arcs -ftest-coverage -fno-omit-frame-pointer -fsanitize-address-use-after-scope -fsanitize=${sanitizers} -fno-sanitize-recover=${sanitizers}"
-        cmake_build("g++-7", "-DCMAKE_BUILD_TYPE=debug ${cmake_linker_flags} -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'")
+        cmake_build("g++-7", "-DCMAKE_BUILD_TYPE=debug -DMIGRAPHX_ENABLE_PYTHON=Off ${cmake_linker_flags} -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'")
 
     }
     stage('Codecov') {

README.md

@@ -8,6 +8,7 @@ AMD's graph optimization engine.
 * [HIP](https://github.com/ROCm-Developer-Tools/HIP) for running on the GPU
 * [Protobuf](https://github.com/google/protobuf) for reading [onxx](https://github.com/onnx/onnx) files
 * [Half](http://half.sourceforge.net/) - IEEE 754-based half-precision floating point library
+* [pybind11](https://pybind11.readthedocs.io/en/stable/) - for python bindings
 
 ## Installing the dependencies
 

requirements.txt

@@ -4,3 +4,4 @@ ROCmSoftwarePlatform/rocBLAS@30a992ae02fda568688bcd190edd5e277d6674d9
 ROCmSoftwarePlatform/MIOpen@1.7.0
 blaze,https://bitbucket.org/blaze-lib/blaze/get/f0755dea0e03.tar.gz -X header -DHEADER_DIR=blaze
 half,https://github.com/pfultz2/half/archive/1.12.0.tar.gz -X header -H sha256:0a08660b68abb176ebc2a0cdf8de46e3182a7f46c66443bb80dbfaaec98cf969
+pybind/pybind11@v2.2.4 -DPYBIND11_TEST=Off --build

src/CMakeLists.txt

@@ -36,6 +36,7 @@ target_include_directories(migraphx SYSTEM PUBLIC $<BUILD_INTERFACE:${HALF_INCLU
 set(PACKAGE_DEPENDS)
 
 add_subdirectory(onnx)
+add_subdirectory(py)
 add_subdirectory(targets/cpu)
 if(MIGRAPHX_ENABLE_GPU)
 list(APPEND PACKAGE_DEPENDS MIOpen rocblas)

src/include/migraphx/program.hpp

@@ -105,6 +105,8 @@ struct program
     void debug_print(instruction_ref ins) const;
     void debug_print(const std::vector<instruction_ref>& inss) const;
 
+    void dry_run(parameter_map params) const;
+
     friend std::ostream& operator<<(std::ostream& os, const program& p);
     friend bool operator==(const program& x, const program& y);
     friend bool operator!=(const program& x, const program& y) { return !(x == y); }

src/include/migraphx/shape.hpp

@@ -62,6 +62,19 @@ struct shape
     shape(type_t t, std::vector<std::size_t> l);
     shape(type_t t, std::vector<std::size_t> l, std::vector<std::size_t> s);
 
+    template <class Range>
+    shape(type_t t, const Range& l) : shape(t, std::vector<std::size_t>(l.begin(), l.end()))
+    {
+    }
+
+    template <class Range1, class Range2>
+    shape(type_t t, const Range1& l, const Range2& s)
+        : shape(t,
+                std::vector<std::size_t>(l.begin(), l.end()),
+                std::vector<std::size_t>(s.begin(), s.end()))
+    {
+    }
+
     type_t type() const;
     const std::vector<std::size_t>& lens() const;
     const std::vector<std::size_t>& strides() const;
@@ -141,6 +154,8 @@ struct shape
         {
             return reinterpret_cast<const T*>(buffer) + n;
         }
+
+        type_t type_enum() const { return get_type<T>{}; }
     };
 
     template <class Visitor>
@@ -156,6 +171,14 @@ struct shape
         MIGRAPHX_THROW("Unknown type");
     }
 
+    template <class Visitor>
+    static void visit_types(Visitor v)
+    {
+#define MIGRAPHX_SHAPE_GENERATE_VISITOR_ALL(x, t) v(as<t>());
+        MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_VISITOR_ALL)
+#undef MIGRAPHX_SHAPE_GENERATE_VISITOR_ALL
+    }
+
     private:
     std::shared_ptr<const shape_impl> impl;
 

src/instruction.cpp

@@ -97,7 +97,7 @@ const std::vector<instruction_ref>& instruction::outputs() const { return output
 
 bool operator==(const instruction& x, const instruction& y)
 {
-    if(not(x.result == y.result and x.op == y.op and x.arguments == y.arguments))
+    if(std::tie(x.result, x.op, x.arguments) != std::tie(y.result, y.op, y.arguments))
         return false;
     if(x.name() == "@literal")
         return x.lit == y.lit;

src/program.cpp

@@ -379,20 +379,31 @@ argument generic_eval(const program& p,
 
 argument program::eval(std::unordered_map<std::string, argument> params) const
 {
+    auto& ctx = this->impl->ctx;
+#ifndef NDEBUG
+    auto sctx          = ctx;
+    auto check_context = [&](auto f) {
+        assert(is_shared(ctx, sctx));
+        auto x = f();
+        sctx   = ctx;
+        return x;
+    };
+#else
+    auto check_context = [](auto f) { return f(); };
+#endif
     if(enabled(MIGRAPHX_TRACE_EVAL{}))
     {
-        auto& ctx = this->impl->ctx;
-        return generic_eval(*this, this->impl->ctx, std::move(params), [&](auto& ins, auto f) {
+        return generic_eval(*this, ctx, std::move(params), [&](auto& ins, auto f) {
             ctx.finish();
             std::cout << "Run instruction: ";
             this->debug_print(ins);
-            return f();
+            return check_context(f);
         });
     }
     else
     {
         return generic_eval(
-            *this, this->impl->ctx, std::move(params), [](auto&, auto f) { return f(); });
+            *this, ctx, std::move(params), [&](auto&, auto f) { return check_context(f); });
     }
 }
 
@@ -446,8 +457,7 @@ void program::perf_report(std::ostream& os, std::size_t n, parameter_map params)
     overhead_vec.reserve(n);
     for(std::size_t i = 0; i < n; i++)
     {
-        overhead_vec.push_back(time<milliseconds>(
-            [&] { generic_eval(*this, ctx, params, [](auto...) { return argument{}; }); }));
+        overhead_vec.push_back(time<milliseconds>([&] { dry_run(params); }));
     }
 
     double total_time             = common_average(total_vec);
@@ -511,6 +521,12 @@ void program::debug_print(const std::vector<instruction_ref>& inss) const
     std::cout << std::endl;
 }
 
+void program::dry_run(std::unordered_map<std::string, argument> params) const
+{
+    auto& ctx = this->impl->ctx;
+    generic_eval(*this, ctx, std::move(params), [](auto&&...) { return argument{}; });
+}
+
 bool operator==(const program& x, const program& y) { return to_string(x) == to_string(y); }
 
 std::ostream& operator<<(std::ostream& os, const program& p)

src/py/CMakeLists.txt

+
+option(MIGRAPHX_ENABLE_PYTHON "Enable python bindings" ON)
+if(MIGRAPHX_ENABLE_PYTHON)
+find_program(DEFAULT_PYTHON_EXE python)
+if(DEFAULT_PYTHON_EXE)
+set(PYTHON_EXECUTABLE ${DEFAULT_PYTHON_EXE} CACHE PATH "Path to python executable")
+endif()
+find_package(pybind11 REQUIRED)
+pybind11_add_module(migraphx_py migraphx_py.cpp)
+set_target_properties(migraphx_py PROPERTIES 
+    OUTPUT_NAME migraphx
+    C_VISIBILITY_PRESET hidden
+    CXX_VISIBILITY_PRESET hidden
+)
+target_link_libraries(migraphx_py PRIVATE migraphx migraphx_onnx migraphx_cpu)
+if(MIGRAPHX_ENABLE_GPU)
+target_link_libraries(migraphx_py PRIVATE migraphx_gpu)
+target_compile_definitions(migraphx_py PRIVATE -DHAVE_GPU)
+endif()
+endif()

src/py/migraphx_py.cpp

+
+#include <pybind11/pybind11.h>
+#include <pybind11/stl.h>
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/cpu/target.hpp>
+#include <migraphx/onnx.hpp>
+#include <migraphx/stringutils.hpp>
+#ifdef HAVE_GPU
+#include <migraphx/gpu/target.hpp>
+#include <migraphx/gpu/hip.hpp>
+#endif
+
+namespace py = pybind11;
+
+template <class F>
+struct throw_half
+{
+    F f;
+
+    template <class A>
+    void operator()(A a) const
+    {
+        f(a);
+    }
+
+    void operator()(migraphx::shape::as<migraphx::half>) const
+    {
+        throw std::runtime_error("Half not supported in python yet.");
+    }
+};
+
+template <class F>
+struct skip_half
+{
+    F f;
+
+    template <class A>
+    void operator()(A a) const
+    {
+        f(a);
+    }
+
+    void operator()(migraphx::shape::as<migraphx::half>) const {}
+};
+
+template <class F>
+void visit_type(const migraphx::shape& s, F f)
+{
+    s.visit_type(throw_half<F>{f});
+}
+
+template <class F>
+void visit_types(F f)
+{
+    migraphx::shape::visit_types(skip_half<F>{f});
+}
+
+template <class T>
+py::buffer_info to_buffer_info(T& x)
+{
+    migraphx::shape s = x.get_shape();
+    py::buffer_info b;
+    visit_type(s, [&](auto as) {
+        b = py::buffer_info(x.data(),
+                            as.size(),
+                            py::format_descriptor<decltype(as())>::format(),
+                            s.lens().size(),
+                            s.lens(),
+                            s.strides());
+    });
+    return b;
+}
+
+migraphx::shape to_shape(const py::buffer_info& info)
+{
+    migraphx::shape::type_t t;
+    visit_types([&](auto as) {
+        if(info.format == py::format_descriptor<decltype(as())>::format())
+            t = as.type_enum();
+    });
+    return migraphx::shape{t, info.shape, info.strides};
+}
+
+PYBIND11_MODULE(migraphx, m)
+{
+    py::class_<migraphx::shape>(m, "shape")
+        .def(py::init<>())
+        .def("type", &migraphx::shape::type)
+        .def("lens", &migraphx::shape::lens)
+        .def("strides", &migraphx::shape::strides)
+        .def("elements", &migraphx::shape::elements)
+        .def("bytes", &migraphx::shape::bytes)
+        .def("type_size", &migraphx::shape::type_size)
+        .def("packed", &migraphx::shape::packed)
+        .def("transposed", &migraphx::shape::transposed)
+        .def("broadcasted", &migraphx::shape::broadcasted)
+        .def("standard", &migraphx::shape::standard)
+        .def("scalar", &migraphx::shape::scalar)
+        .def("__eq__", std::equal_to<migraphx::shape>{})
+        .def("__ne__", std::not_equal_to<migraphx::shape>{})
+        .def("__repr__", [](const migraphx::shape& s) { return migraphx::to_string(s); });
+
+    py::class_<migraphx::argument>(m, "argument", py::buffer_protocol())
+        .def_buffer([](migraphx::argument& x) -> py::buffer_info { return to_buffer_info(x); })
+        .def("__init__",
+             [](migraphx::argument& x, py::buffer b) {
+                 py::buffer_info info = b.request();
+                 new(&x) migraphx::argument(to_shape(info), info.ptr);
+             })
+        .def("__eq__", std::equal_to<migraphx::argument>{})
+        .def("__ne__", std::not_equal_to<migraphx::argument>{})
+        .def("__repr__", [](const migraphx::argument& x) { return migraphx::to_string(x); });
+
+    py::class_<migraphx::target>(m, "target");
+
+    py::class_<migraphx::program>(m, "program")
+        .def("get_parameter_shapes", &migraphx::program::get_parameter_shapes)
+        .def("get_shape", &migraphx::program::get_shape)
+        .def("compile", [](migraphx::program& p, const migraphx::target& t) { p.compile(t); })
+        .def("run", &migraphx::program::eval)
+        .def("__eq__", std::equal_to<migraphx::program>{})
+        .def("__ne__", std::not_equal_to<migraphx::program>{})
+        .def("__repr__", [](const migraphx::program& p) { return migraphx::to_string(p); });
+
+    m.def("parse_onnx", &migraphx::parse_onnx);
+
+    m.def("get_target", [](const std::string& name) -> migraphx::target {
+        if(name == "cpu")
+            return migraphx::cpu::target{};
+#ifdef HAVE_GPU
+        if(name == "gpu")
+            return migraphx::gpu::target{};
+#endif
+        throw std::runtime_error("Target not found: " + name);
+    });
+
+    m.def("generate_argument", &migraphx::generate_argument, py::arg("s"), py::arg("seed") = 0);
+
+#ifdef HAVE_GPU
+    m.def("allocate_gpu", &migraphx::gpu::allocate_gpu, py::arg("s"), py::arg("host") = false);
+    m.def("to_gpu", &migraphx::gpu::to_gpu, py::arg("arg"), py::arg("host") = false);
+    m.def("from_gpu", &migraphx::gpu::from_gpu);
+    m.def("gpu_sync", &migraphx::gpu::gpu_sync);
+    m.def("copy_to_gpu", &migraphx::gpu::copy_to_gpu);
+#endif
+
+#ifdef VERSION_INFO
+    m.attr("__version__") = VERSION_INFO;
+#else
+    m.attr("__version__") = "dev";
+#endif
+}

test/CMakeLists.txt

@@ -112,7 +112,8 @@ if(MIGRAPHX_ENABLE_GPU)
 endif()
 
 # Onnx test
-add_executable(test_onnx onnx/onnx_test.cpp)
+set(TEST_ONNX_DIR ${CMAKE_CURRENT_SOURCE_DIR}/onnx)
+add_executable(test_onnx ${TEST_ONNX_DIR}/onnx_test.cpp)
 rocm_clang_tidy_check(test_onnx)
 target_link_libraries(test_onnx migraphx_onnx)
 target_include_directories(test_onnx PUBLIC include)
@@ -120,6 +121,10 @@ add_test(NAME test_onnx COMMAND $<TARGET_FILE:test_onnx> WORKING_DIRECTORY ${CMA
 add_dependencies(tests test_onnx)
 add_dependencies(check test_onnx)
 
+if(MIGRAPHX_ENABLE_PYTHON)
+add_subdirectory(py)
+endif()
+
 function(test_header NAME HEADER)
 
     file(WRITE ${CMAKE_CURRENT_BINARY_DIR}/header-main-include-${NAME}.cpp 

test/gpu/miopen.cpp

@@ -122,7 +122,14 @@ migraphx::argument run_gpu(migraphx::program& p)
         m[x.first] =
             migraphx::gpu::to_gpu(migraphx::generate_argument(x.second, get_hash(x.first)));
     }
+    // Program should have an output parameter
     EXPECT(bool{m.find("output") != m.end()});
+    // Ensure the program doesn't modify the context in a dry run
+    auto ctx = p.get_context();
+    assert(&ctx != &p.get_context());
+    EXPECT(is_shared(ctx, p.get_context()));
+    p.dry_run(m);
+    EXPECT(is_shared(ctx, p.get_context()));
     return migraphx::gpu::from_gpu(p.eval(m));
 }
 

test/py/CMakeLists.txt

+
+find_package(PythonInterp)
+
+function(add_py_test NAME SCRIPT)
+    set (ENV_COMMAND ${CMAKE_COMMAND} -E env 
+        "PYTHONPATH=$<TARGET_FILE_DIR:migraphx_py>"
+    )
+    add_test(
+        NAME test_py_${NAME} 
+        COMMAND ${ENV_COMMAND} ${PYTHON_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/${SCRIPT} ${ARGN})
+    add_custom_target(test_py_${NAME} 
+        COMMAND ${ENV_COMMAND} ${PYTHON_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/${SCRIPT} ${ARGN} 
+        COMMENT "${PYTHON_EXECUTABLE} ${SCRIPT}")
+endfunction()
+add_dependencies(tests migraphx_py)
+add_dependencies(check migraphx_py)
+
+add_py_test(cpu cpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
+if(MIGRAPHX_ENABLE_GPU)
+add_py_test(gpu gpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
+endif()

test/py/cpu.py

+import migraphx
+
+p = migraphx.parse_onnx("conv_relu_maxpool.onnx")
+print(p)
+s1 = p.get_shape()
+print("Compiling ...")
+p.compile(migraphx.get_target("cpu"))
+print(p)
+s2 = p.get_shape()
+assert s1 == s2
+params = {}
+for key, value in p.get_parameter_shapes().items():
+    print("Parameter {} -> {}".format(key, value))
+    params[key] = migraphx.generate_argument(value)
+
+r = p.run(params)
+print(r)

test/py/gpu.py

+import migraphx
+
+p = migraphx.parse_onnx("conv_relu_maxpool.onnx")
+print(p)
+print("Compiling ...")
+p.compile(migraphx.get_target("gpu"))
+print(p)
+params = {}
+for key, value in p.get_parameter_shapes().items():
+    print("Parameter {} -> {}".format(key, value))
+    params[key] = migraphx.to_gpu(migraphx.generate_argument(value))
+
+r = migraphx.from_gpu(p.run(params))
+print(r)