Add test for miopen

.clang-tidy

@@ -97,7 +97,7 @@ CheckOptions:
     value:           CamelCase
   - key:             readability-identifier-naming.TypeAliasCase
     value:           lower_case
-  - key:             readability-identifier-naming.MacroDefinitionCase
-    value:           UPPER_CASE
-  - key:             readability-identifier-naming.MacroDefinitionPrefix
-    value:           RTG_
+  # - key:             readability-identifier-naming.MacroDefinitionCase
+  #   value:           UPPER_CASE
+  # - key:             readability-identifier-naming.MacroDefinitionPrefix
+  #   value:           RTG_

CMakeLists.txt

@@ -14,15 +14,15 @@ add_compile_options(-std=c++14)
 list(APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake)
 include(EnableCompilerWarnings)
 # Override clang-tidy to not find the version from hcc
-find_program(CLANG_TIDY_EXE 
-    NAMES 
-        clang-tidy
-        clang-tidy-5.0
-        clang-tidy-6.0
-        clang-tidy-7.0
-    PATHS
-        /usr/local/opt/llvm/bin
-)
+# find_program(CLANG_TIDY_EXE 
+#     NAMES 
+#         clang-tidy
+#         clang-tidy-5.0
+#         clang-tidy-6.0
+#         clang-tidy-7.0
+#     PATHS
+#         /usr/local/opt/llvm/bin
+# )
 include(ROCMClangTidy)
 rocm_enable_clang_tidy(
     CHECKS
@@ -52,10 +52,12 @@ rocm_enable_clang_tidy(
         -hicpp-explicit-conversions
         -hicpp-no-array-decay
         -hicpp-special-member-functions
+        -hicpp-use-override
         -llvm-header-guard
         -llvm-include-order
         -misc-macro-parentheses
         -modernize-use-auto
+        -modernize-use-override
         -modernize-pass-by-value
         -modernize-use-default-member-init
         -modernize-use-transparent-functors

dev-requirements.txt

+pfultz2/rocm-recipes
+-f requirements.txt

src/include/rtg/argument.hpp

@@ -26,6 +26,8 @@ struct argument : raw_data<argument>
     }
 
     argument(shape s, std::function<char*()> d) : data(d), m_shape(s) {}
+    template<class T>
+    argument(shape s, T* d) : data([d] { return reinterpret_cast<char*>(d); }), m_shape(s) {}
 
     /// Provides a raw pointer to the data
     std::function<char*()> data;

src/include/rtg/builtin.hpp

@@ -15,6 +15,14 @@ struct literal
     argument compute(shape, std::vector<argument>) const { RTG_THROW("builtin"); }
 };
 
+struct outline
+{
+    shape s;
+    std::string name() const { return "@outline"; }
+    shape compute_shape(std::vector<shape>) const { RTG_THROW("builtin"); }
+    argument compute(shape, std::vector<argument>) const { RTG_THROW("builtin"); }
+};
+
 struct param
 {
     std::string parameter;

src/include/rtg/manage_ptr.hpp

@@ -6,7 +6,7 @@
 
 namespace rtg {
 
-template <class F, F f>
+template <class F, F f> // NOLINT
 struct manage_deleter
 {
     template <class T>
@@ -43,6 +43,12 @@ using remove_ptr = typename std::
 template <class T>
 using shared = std::shared_ptr<remove_ptr<T>>;
 
+template<class T>
+shared<T> share(T p)
+{
+    return shared<T>{std::move(p)};
+}
+
 } // namespace rtg
 
 #define RTG_MANAGE_PTR(T, F) rtg::manage_ptr<std::remove_pointer_t<T>, decltype(&F), &F> // NOLINT

src/include/rtg/operators.hpp

@@ -218,6 +218,21 @@ struct reshape
     }
 };
 
+struct outline
+{
+    shape s;
+    std::string name() const { return "outline"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(0);
+        return s;
+    }
+    argument compute(shape, std::vector<argument>) const
+    {
+        return {s, nullptr};
+    }
+};
+
 } // namespace rtg
 
 #endif

src/include/rtg/program.hpp

@@ -56,8 +56,12 @@ struct program
 
     instruction_ref add_literal(literal l);
     
+    instruction_ref add_outline(shape s);
+
     instruction_ref add_parameter(std::string name, shape s);
 
+    shape get_parameter_shape(std::string name);
+
     literal eval(std::unordered_map<std::string, argument> params) const;
 
     friend std::ostream& operator<<(std::ostream& os, const program& p);

src/program.cpp

@@ -51,14 +51,37 @@ program::replace_instruction(instruction_ref ins, operation op, std::vector<inst
 
 instruction_ref program::add_literal(literal l)
 {
-    impl->instructions.emplace_back(std::move(l));
-    return std::prev(impl->instructions.end());
+    impl->instructions.emplace_front(std::move(l));
+    return impl->instructions.begin();
+}
+
+instruction_ref program::add_outline(shape s)
+{
+    impl->instructions.push_front({builtin::outline{s}, s, {}});
+    return impl->instructions.begin();
 }
 
 instruction_ref program::add_parameter(std::string name, shape s)
 {
-    impl->instructions.push_back({builtin::param{std::move(name)}, s, {}});
-    return std::prev(impl->instructions.end());
+    impl->instructions.push_front({builtin::param{std::move(name)}, s, {}});
+    return impl->instructions.begin();
+}
+
+shape program::get_parameter_shape(std::string name)
+{
+    auto ins = std::find_if(
+               impl->instructions.begin(), impl->instructions.end(), [&](const instruction& x) {
+                   if(x.op.name() == "@param") 
+                   {
+                        return any_cast<builtin::param>(x.op).parameter == name;
+                   }
+                   else
+                   {
+                        return false;
+                   }
+               });
+    if (ins != this->end()) return ins->result;
+    else return {};
 }
 
 bool program::has_instruction(instruction_ref ins) const
@@ -102,6 +125,10 @@ literal program::eval(std::unordered_map<std::string, argument> params) const
         {
             result = params.at(any_cast<builtin::param>(ins.op).parameter);
         }
+        else if(ins.op.name() == "@outline")
+        {
+            result = argument{ins.result, nullptr};
+        }
         else
         {
             std::vector<argument> values(ins.arguments.size());

src/targets/cpu/cpu_target.cpp

@@ -26,7 +26,7 @@ struct cpu_convolution
             auto wei_h = weights.get_shape().lens()[2];
             auto wei_w = weights.get_shape().lens()[3];
 
-            dfor(in_n, in_c, in_h, in_w)(
+            dfor(output_shape.lens()[0], output_shape.lens()[1], output_shape.lens()[2], output_shape.lens()[3])(
                 [&](std::size_t o, std::size_t w, std::size_t i, std::size_t j) {
                     const int start_x = i * op.stride[0] - op.padding[0];
                     const int start_y = j * op.stride[1] - op.padding[1];

src/targets/miopen/miopen_target.cpp

@@ -8,6 +8,41 @@
 namespace rtg {
 namespace miopen {
 
+
+struct hip_allocate
+{
+    std::string name() const { return "hip::allocate"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return inputs.front();
+    }
+    argument compute(shape output_shape, std::vector<argument>) const
+    {
+        char * data = nullptr;
+        // TODO: Check return status
+        hipMalloc(&data, output_shape.bytes());
+        return {output_shape, data};
+    }
+};
+
+struct hip_free
+{
+    std::string name() const { return "hip::free"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return {};
+    }
+    argument compute(shape output_shape, std::vector<argument> args) const
+    {
+        // TODO: Check return status
+        hipFree(args.front().data());
+        return {};
+    }
+};
+
+
 using miopen_handle     = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
 using tensor_descriptor = RTG_MANAGE_PTR(miopenTensorDescriptor_t, miopenDestroyTensorDescriptor);
 using convolution_descriptor = RTG_MANAGE_PTR(miopenConvolutionDescriptor_t,
@@ -55,19 +90,26 @@ convolution_descriptor make_conv(const rtg::convolution& op)
     return c;
 }
 
+activation_descriptor make_relu()
+{
+    auto ad = make_obj<activation_descriptor>(&miopenCreateActivationDescriptor);
+    miopenSetActivationDescriptor(ad.get(), miopenActivationRELU, 0, 0, 0);
+    return ad;
+}
+
 struct miopen_convolution
 {
     convolution op;
-    convolution_descriptor cd;
+    shared<convolution_descriptor> cd;
 
     std::string name() const { return "miopen::convolution"; }
     shape compute_shape(std::vector<shape> inputs) const
     {
+        check_shapes{inputs}.has(4);
         return op.compute_shape({inputs.at(1), inputs.at(2)});
     }
     argument compute(shape output_shape, std::vector<argument> args) const
     {
-        argument result;
         auto x_desc = make_tensor(args[1].get_shape());
         auto w_desc = make_tensor(args[2].get_shape());
         auto y_desc = make_tensor(output_shape);
@@ -82,12 +124,12 @@ struct miopen_convolution
                                               args[2].get(),
                                               cd.get(),
                                               y_desc.get(),
-                                              args[4].get(),
+                                              args[3].get(),
                                               1,
                                               &algo_count,
                                               &perf,
-                                              args[3].get(),
-                                              args[3].get_shape().bytes(),
+                                              nullptr,
+                                              0,
                                               false);
         miopenConvolutionForward(args[0].get(),
                                  &alpha,
@@ -99,19 +141,42 @@ struct miopen_convolution
                                  perf.fwd_algo,
                                  &beta,
                                  y_desc.get(),
-                                 args[4].get(),
                                  args[3].get(),
-                                 args[3].get_shape().bytes());
-        return result;
+                                 nullptr,
+                                 0);
+        return args[3];
+    }
+};
+
+struct miopen_relu
+{
+    shared<activation_descriptor> ad;
+    std::string name() const { return "miopen::relu"; }
+    shape compute_shape(std::vector<shape> inputs) const 
+    {
+        check_shapes{inputs}.has(3); 
+        return inputs.at(1); 
+    }
+
+    argument compute(shape output_shape, std::vector<argument> args) const
+    {
+        float alpha = 1, beta = 0;
+        auto x_desc = make_tensor(args[1].get_shape());
+        auto y_desc = make_tensor(output_shape);
+        miopenActivationForward(args[0].get(), ad.get(), &alpha, x_desc.get(), args[1].get(), &beta, y_desc.get(), args[2].get());
+
+        return args[2];
     }
 };
 
 struct miopen_apply
 {
     program* prog;
+    instruction_ref handle;
 
     void apply()
     {
+        handle = prog->add_parameter("handle", shape{shape::any_type});
         for(auto it = prog->begin(); it != prog->end(); it++)
         {
             if(it->op.name() == "convolution")
@@ -125,13 +190,40 @@ struct miopen_apply
         }
     }
 
+    instruction_ref insert_allocation(instruction_ref ins, const shape& s)
+    {
+        if (ins == --prog->end())
+        {
+            return prog->add_parameter("output", s);
+        }
+        else
+        {
+            auto is = prog->add_outline(s);
+            auto result = prog->insert_instruction(ins, hip_allocate{}, is);
+            prog->insert_instruction(++ins, hip_free{}, result);
+            return result;
+        }
+    }
+
     void apply_convolution(instruction_ref ins)
     {
-        // auto&& op = any_cast<convolution>(ins->op);
-        // prog->replace_instruction(ins, miopen_convolution{op}, ins->arguments);
+        auto&& op = any_cast<convolution>(ins->op);
+        auto cd = make_conv(op);
+        auto output = insert_allocation(ins, ins->result);
+
+        prog->replace_instruction(ins, miopen_convolution{op, std::move(cd)}, handle, ins->arguments.at(0), ins->arguments.at(1), output);
     }
 
-    void apply_activation(instruction_ref ins) {}
+    void apply_activation(instruction_ref ins) 
+    {
+        auto&& op = any_cast<activation>(ins->op);
+        auto ad = make_relu();
+        if(op.mode == "relu") 
+        {
+            auto output = insert_allocation(ins, ins->result);
+            prog->replace_instruction(ins, miopen_relu{std::move(ad)}, handle, ins->arguments.at(0), output);
+        }
+    }
 };
 
 std::string miopen_target::name() const { return "miopen"; }

test/CMakeLists.txt

@@ -59,7 +59,7 @@ function(add_test_executable TEST_NAME)
     add_dependencies(tests ${TEST_NAME})
     add_dependencies(check ${TEST_NAME})
     set_tests_properties(${TEST_NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED")
-    target_link_libraries(${TEST_NAME} rtg)
+    target_link_libraries(${TEST_NAME} rtg rtg_cpu rtg_miopen)
 endfunction(add_test_executable)
 
 file(GLOB TESTS *.cpp)

test/miopen.cpp

+
+#include <rtg/program.hpp>
+#include <rtg/operators.hpp>
+#include <rtg/cpu/cpu_target.hpp>
+#include <rtg/miopen/miopen_target.hpp>
+#include <rtg/manage_ptr.hpp>
+
+#include <miopen/miopen.h>
+
+#include <random>
+
+#include "test.hpp"
+
+using hip_ptr = RTG_MANAGE_PTR(void, hipFree);
+using miopen_handle     = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
+
+template <class Result, class F, class... Ts>
+Result make_obj(F f, Ts... xs)
+{
+    typename Result::pointer x = nullptr;
+    auto status                = f(&x, xs...);
+    Result r{x};
+    if(status != miopenStatusSuccess)
+        RTG_THROW("MIOpen call failed");
+    return r;
+}
+
+hip_ptr hip_allocate(std::size_t sz)
+{
+    void* result;
+    // TODO: Check status
+    hipMalloc(&result, sz);
+    return hip_ptr{result};
+}
+
+template<class T>
+hip_ptr write(const T& x)
+{
+    using type = typename T::value_type;
+    auto size = x.size() * sizeof(type);
+    auto result = hip_allocate(size);
+    // TODO: Check status
+    hipMemcpy(result.get(), x.data(), size, hipMemcpyHostToDevice);
+    return result;
+}
+
+template<class T>
+std::vector<T> read(const hip_ptr& x, std::size_t sz)
+{
+    std::vector<T> result(sz);
+    // TODO: Check status
+    hipMemcpy(result.data(), x.get(), sz * sizeof(T), hipMemcpyDeviceToHost);
+    return result;
+}
+
+rtg::program create_program()
+{
+    rtg::program p;
+    auto input = p.add_parameter("x", rtg::shape{rtg::shape::float_type, {4, 3, 3, 3}});
+    auto weights = p.add_parameter("w", rtg::shape{rtg::shape::float_type, {4, 3, 3, 3}});
+    auto conv = p.add_instruction(rtg::convolution{}, input, weights);
+    p.add_instruction(rtg::activation{"relu"}, conv);
+    return p;
+}
+
+std::vector<float> get_tensor_data(rtg::shape s)
+{
+    std::vector<float> result(s.elements());
+    std::mt19937 engine{0};
+    std::uniform_real_distribution<> dist;
+    std::generate(result.begin(), result.end(), [&] { return dist(engine); });
+    return result;
+}
+
+rtg::argument get_tensor_argument_cpu(rtg::shape s)
+{
+    auto v = get_tensor_data(s);
+    return {s, [v]() mutable { return reinterpret_cast<char*>(v.data()); }};
+}
+
+rtg::argument get_tensor_argument_gpu(rtg::shape s)
+{
+    auto v = get_tensor_data(s);
+    auto p = rtg::share(write(v));
+    return {s, [p]() mutable { return reinterpret_cast<char*>(p.get()); }};
+}
+
+std::vector<float> cpu()
+{
+    std::vector<float> result;
+    auto p = create_program();
+    auto x = get_tensor_argument_cpu({rtg::shape::float_type, {4, 3, 3, 3}});
+    auto w = get_tensor_argument_cpu({rtg::shape::float_type, {4, 3, 3, 3}});
+    p.compile(rtg::cpu::cpu_target{});
+    auto r = p.eval({
+        {"x", x}, 
+        {"w", w}
+    });
+    r.visit([&](auto output) { result.assign(output.begin(), output.end()); });
+    return result;
+}
+
+std::vector<float> gpu()
+{
+    std::vector<float> result;
+    auto p = create_program();
+    auto x = get_tensor_argument_gpu({rtg::shape::float_type, {4, 3, 3, 3}});
+    auto w = get_tensor_argument_gpu({rtg::shape::float_type, {4, 3, 3, 3}});
+    p.compile(rtg::miopen::miopen_target{});
+    auto y = get_tensor_argument_gpu(p.get_parameter_shape("output"));
+    auto handle = make_obj<miopen_handle>(&miopenCreate);
+    auto r = p.eval({
+        {"x", x}, 
+        {"w", w},
+        {"output", y},
+        {"handle", {rtg::shape::any_type, handle.get()}}
+    });
+    r.visit([&](auto output) { result.assign(output.begin(), output.end()); });
+    return result;
+}
+
+void test1() 
+{
+    auto x = cpu();
+    auto y = gpu();
+    if (x == y)
+        printf("FAILED\n");
+}
+
+int main()
+{
+    test1();
+}

test/test.hpp

@@ -82,7 +82,7 @@ struct lhs_expression
     template <class U>                            \
     auto operator op(const U& rhs) const          \
     {                                             \
-        return make_expression(lhs, rhs, name{}); \
+        return make_expression(lhs, rhs, name{}); /* NOLINT */ \
     }
 
     TEST_FOREACH_OPERATOR(TEST_LHS_OPERATOR)