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

CMakeLists.txt

@@ -145,6 +145,7 @@ rocm_create_package(
     DESCRIPTION "AMD's graph optimizer"
     MAINTAINER "Paul Fultz II <paul.fultz@amd.com>"
     LDCONFIG
+    PTH
     DEPENDS miopen-hip rocblas hip_hcc half
 )
 

Jenkinsfile

@@ -107,6 +107,10 @@ rocmtest tidy: rocmnode('rocmtest') { cmake_build ->
     stage('Clang Release') {
         cmake_build("hcc", "-DCMAKE_BUILD_TYPE=release")
     }
+
+    stage('Clang Release Python 3') {
+        cmake_build("hcc", "-DCMAKE_BUILD_TYPE=release -DPYTHON_EXECUTABLE=/usr/local/bin/python3")
+    }
 }, gcc5: rocmnode('rocmtest') { cmake_build ->
     stage('GCC 5 Debug') {
         cmake_build("g++-5", "-DCMAKE_BUILD_TYPE=debug")

requirements.txt

 google/protobuf -DCMAKE_POSITION_INDEPENDENT_CODE=On
-RadeonOpenCompute/rocm-cmake@ac45c6e --build
+RadeonOpenCompute/rocm-cmake@42f6740 --build
 ROCmSoftwarePlatform/rocBLAS@30a992ae02fda568688bcd190edd5e277d6674d9
 ROCmSoftwarePlatform/MIOpen@1.7.0
 blaze,https://bitbucket.org/blaze-lib/blaze/get/f0755dea0e03.tar.gz -X header -DHEADER_DIR=blaze

src/include/migraphx/tensor_view.hpp

@@ -124,6 +124,8 @@ struct tensor_view
             return m_data + this->size();
     }
 
+    std::vector<T> to_vector() const { return std::vector<T>(this->begin(), this->end()); }
+
     friend std::ostream& operator<<(std::ostream& os, const tensor_view<T>& x)
     {
         if(!x.empty())

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()
+    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()
+    rocm_install_targets(TARGETS migraphx_py)
 endif()

src/py/migraphx_py.cpp

@@ -28,6 +28,11 @@ struct throw_half
     {
         throw std::runtime_error("Half not supported in python yet.");
     }
+
+    void operator()(migraphx::tensor_view<migraphx::half>) const
+    {
+        throw std::runtime_error("Half not supported in python yet.");
+    }
 };
 
 template <class F>
@@ -42,6 +47,8 @@ struct skip_half
     }
 
     void operator()(migraphx::shape::as<migraphx::half>) const {}
+
+    void operator()(migraphx::tensor_view<migraphx::half>) const {}
 };
 
 template <class F>
@@ -50,6 +57,12 @@ void visit_type(const migraphx::shape& s, F f)
     s.visit_type(throw_half<F>{f});
 }
 
+template <class T, class F>
+void visit(const migraphx::raw_data<T>& x, F f)
+{
+    x.visit(throw_half<F>{f});
+}
+
 template <class F>
 void visit_types(F f)
 {
@@ -60,6 +73,9 @@ template <class T>
 py::buffer_info to_buffer_info(T& x)
 {
     migraphx::shape s = x.get_shape();
+    auto strides      = s.strides();
+    std::transform(
+        strides.begin(), strides.end(), strides.begin(), [&](auto i) { return i * s.type_size(); });
     py::buffer_info b;
     visit_type(s, [&](auto as) {
         b = py::buffer_info(x.data(),
@@ -67,7 +83,7 @@ py::buffer_info to_buffer_info(T& x)
                             py::format_descriptor<decltype(as())>::format(),
                             s.lens().size(),
                             s.lens(),
-                            s.strides());
+                            strides);
     });
     return b;
 }
@@ -75,11 +91,20 @@ py::buffer_info to_buffer_info(T& x)
 migraphx::shape to_shape(const py::buffer_info& info)
 {
     migraphx::shape::type_t t;
+    std::size_t n = 0;
     visit_types([&](auto as) {
         if(info.format == py::format_descriptor<decltype(as())>::format())
+        {
             t = as.type_enum();
+            n = sizeof(as());
+        }
+
     });
-    return migraphx::shape{t, info.shape, info.strides};
+    auto strides = info.strides;
+    std::transform(strides.begin(), strides.end(), strides.begin(), [&](auto i) -> std::size_t {
+        return n > 0 ? i / n : 0;
+    });
+    return migraphx::shape{t, info.shape, strides};
 }
 
 PYBIND11_MODULE(migraphx, m)
@@ -108,6 +133,13 @@ PYBIND11_MODULE(migraphx, m)
                  py::buffer_info info = b.request();
                  new(&x) migraphx::argument(to_shape(info), info.ptr);
              })
+        .def("get_shape", &migraphx::argument::get_shape)
+        .def("tolist",
+             [](migraphx::argument& x) {
+                 py::list l{x.get_shape().elements()};
+                 visit(x, [&](auto data) { l = py::cast(data.to_vector()); });
+                 return l;
+             })
         .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); });

test/py/CMakeLists.txt

@@ -4,6 +4,8 @@ find_package(PythonInterp)
 function(add_py_test NAME SCRIPT)
     set (ENV_COMMAND ${CMAKE_COMMAND} -E env 
         "PYTHONPATH=$<TARGET_FILE_DIR:migraphx_py>"
+        "PYTHONMALLOC=debug"
+        "MALLOC_CHECK_=3"
     )
     add_test(
         NAME test_py_${NAME} 
@@ -15,7 +17,8 @@ endfunction()
 add_dependencies(tests migraphx_py)
 add_dependencies(check migraphx_py)
 
-add_py_test(cpu cpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
+add_py_test(cpu test_cpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
 if(MIGRAPHX_ENABLE_GPU)
-add_py_test(gpu gpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
+add_py_test(gpu test_gpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
+add_py_test(array test_array.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
 endif()

test/py/test_array.py

+import migraphx, struct, array, sys
+try:
+    from functools import reduce
+except:
+    pass
+
+def assert_eq(x, y):
+    if x == y:
+        pass
+    else:
+        raise Exception(str(x) + " != " + str(y))
+
+def read_float(b, index):
+    return struct.unpack_from('f', b, index*4)[0]
+
+def write_float(b, index):
+    struct.pack_into('f', b, index*4)
+
+def nelements(lens):
+    return reduce(lambda x,y: x*y,lens, 1)
+
+def create_buffer(t, data, shape):
+    a = array.array(t, data)
+    if sys.version_info >= (3, 0):
+        m = memoryview(a.tobytes())
+        return m.cast(t, shape)
+    else:
+        m = memoryview(a.tostring())
+        return m
+
+def check_argument(a):
+    l = a.tolist()
+    for i in range(len(l)):
+        assert_eq(l[i], read_float(a, i))
+
+def check_shapes(r, m):
+    lens = list(m.shape)
+    strides = [int(s/m.itemsize) for s in m.strides]
+    elements = nelements(lens)
+    assert_eq(r.get_shape().elements(), elements)
+    assert_eq(r.get_shape().lens(), lens)
+    assert_eq(r.get_shape().strides(), strides)
+
+def run(p):
+    params = {}
+    for key, value in p.get_parameter_shapes().items():
+        params[key] = migraphx.to_gpu(migraphx.generate_argument(value))
+
+    return migraphx.from_gpu(p.run(params))
+
+def test_shape(shape):
+    data = list(range(nelements(shape)))
+    m = create_buffer('f', data, shape)
+    a = migraphx.argument(m)
+    check_shapes(a, m)
+    assert_eq(a.tolist(), data)
+
+def test_input():
+    if sys.version_info >= (3, 0):
+        test_shape([4])
+        test_shape([2, 3])
+    else:
+        data = list(range(4))
+        m = create_buffer('f', data, [4])
+        a1 = migraphx.argument(m)
+        a2 = migraphx.argument(bytearray(a1))
+        check_shapes(a2, m)
+        assert_eq(a1.tolist(), m.tolist())
+
+
+def test_output():
+    p = migraphx.parse_onnx("conv_relu_maxpool.onnx")
+    p.compile(migraphx.get_target("gpu"))
+
+    r1 = run(p)
+    r2 = run(p)
+    assert_eq(r1, r2)
+    assert_eq(r1.tolist(), r2.tolist())
+
+    check_argument(r1)
+    check_argument(r2)
+
+    m1 = memoryview(r1)
+    m2 = memoryview(r2)
+
+    check_shapes(r1, m1)
+    check_shapes(r2, m2)
+
+
+test_input()
+test_output()