Added more convolution tests

.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

@@ -9,20 +9,27 @@ if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
     endif()
 endif()
 
+if(CMAKE_CXX_COMPILER MATCHES ".*hcc")
+    message(STATUS "Enable miopen backend")
+    set(RTG_ENABLE_MIOPEN On CACHE BOOL "")
+else()
+    set(RTG_ENABLE_MIOPEN Off CACHE BOOL "")
+endif()
+
 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 +59,14 @@ rocm_enable_clang_tidy(
         -hicpp-explicit-conversions
         -hicpp-no-array-decay
         -hicpp-special-member-functions
+        -hicpp-use-override
+        # This check is broken
+        -hicpp-use-auto
         -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
@@ -89,11 +100,12 @@ rocm_enable_cppcheck(
         unusedStructMember
     FORCE
     SOURCES
-        include/
         src/
         test/
     INCLUDE
-        ${CMAKE_CURRENT_SOURCE_DIR}/include
+        ${CMAKE_CURRENT_SOURCE_DIR}/src/include
+        ${CMAKE_CURRENT_SOURCE_DIR}/src/targets/cpu/include
+        ${CMAKE_CURRENT_SOURCE_DIR}/src/targets/miopen/include
 )
 
 add_subdirectory(src)

Dockerfile

@@ -7,7 +7,7 @@ RUN dpkg --add-architecture i386
 
 # Add rocm repository
 RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y curl apt-utils wget
-RUN curl https://raw.githubusercontent.com/RadeonOpenCompute/ROCm-docker/develop/add-rocm.sh | bash
+RUN curl https://raw.githubusercontent.com/RadeonOpenCompute/ROCm-docker/master/add-rocm.sh | bash
 
 # Install dependencies
 RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-unauthenticated \
@@ -19,8 +19,10 @@ RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-
     cmake \
     curl \
     doxygen \
+    gdb \
     git \
     hcc \
+    hip_hcc \
     lcov \
     libnuma-dev \
     python \
@@ -36,14 +38,15 @@ RUN apt-get update && DEBIAN_FRONTEND=noninteractive apt-get install -y --allow-
 # Install cget
 RUN pip install cget
 
-# Install cppcheck
-RUN cget -p $PREFIX install danmar/cppcheck@ab02595be1b17035b534db655f9e119080a368bc
-
-RUN cget -p $PREFIX install pfultz2/rocm-recipes
+# Use hcc
+RUN cget -p $PREFIX init --cxx /opt/rocm/bin/hcc
 
 # Install dependencies
+ADD dev-requirements.txt /dev-requirements.txt
 ADD requirements.txt /requirements.txt
-RUN cget -p $PREFIX install -f /requirements.txt
+RUN cget -p $PREFIX install -f /dev-requirements.txt
+
+ENV LD_LIBRARY_PATH=$PREFIX/lib
 
 # Install doc requirements
 # ADD doc/requirements.txt /doc-requirements.txt

Jenkinsfile

@@ -25,7 +25,7 @@ def rocmtestnode(variant, name, body) {
 
             }
         }
-        withDockerContainer(image: image, args: '--device=/dev/kfd --device=/dev/dri --group-add video') {
+        withDockerContainer(image: image, args: '--device=/dev/kfd --device=/dev/dri --group-add video --cap-add SYS_PTRACE') {
             timeout(time: 1, unit: 'HOURS') {
                 body(cmake_build)
             }
@@ -92,10 +92,10 @@ rocmtest tidy: rocmnode('rocmtest') { cmake_build ->
     }
 }, clang: rocmnode('rocmtest') { cmake_build ->
     stage('Clang Debug') {
-        cmake_build('clang++-5.0', '-DCMAKE_BUILD_TYPE=debug')
+        cmake_build('/opt/rocm/bin/hcc', '-DCMAKE_BUILD_TYPE=debug')
     }
     stage('Clang Release') {
-        cmake_build('clang++-5.0', '-DCMAKE_BUILD_TYPE=release')
+        cmake_build('/opt/rocm/bin/hcc', '-DCMAKE_BUILD_TYPE=release')
     }
 }, gcc: rocmnode('rocmtest') { cmake_build ->
     stage('GCC Debug') {

dev-requirements.txt

+pfultz2/rocm-recipes
+danmar/cppcheck@d9f9bdda7344e80585f71141be7797055d7987f3
+-f requirements.txt

install_deps.cmake

+#!/usr/bin/cmake -P
+
+set(ARGS)
+foreach(i RANGE 3 ${CMAKE_ARGC})
+    list(APPEND ARGS ${CMAKE_ARGV${i}})
+endforeach()
+
+include(CMakeParseArguments)
+
+set(options help)
+set(oneValueArgs --prefix)
+set(multiValueArgs)
+
+cmake_parse_arguments(PARSE "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGS})
+
+if(PARSE_help)
+message("Usage: install_deps.cmake [options] [cmake-args]")
+message("")
+message("Options:")
+message("  --prefix               Set the prefix to install the dependencies.")
+message("")
+message("Commands:")
+message("  help                   Show this message and exit.")
+message("")
+return()
+endif()
+
+set(_PREFIX /usr/local)
+if(PARSE_--prefix)
+    set(_PREFIX ${PARSE_--prefix})
+endif()
+
+get_filename_component(PREFIX ${_PREFIX} ABSOLUTE)
+
+find_package(CMakeGet QUIET PATHS ${PREFIX})
+
+if(NOT CMakeGet_FOUND)
+    set(FILENAME ${PREFIX}/tmp/cmake-get-install.cmake)
+    file(DOWNLOAD https://raw.githubusercontent.com/pfultz2/cmake-get/master/install.cmake ${FILENAME} STATUS RESULT_LIST)
+    list(GET RESULT_LIST 0 RESULT)
+    list(GET RESULT_LIST 1 RESULT_MESSAGE)
+    if(NOT RESULT EQUAL 0)
+        message(FATAL_ERROR "Download for install.cmake failed: ${RESULT_MESSAGE}")
+    endif()
+    execute_process(COMMAND ${CMAKE_COMMAND} -P ${FILENAME} ${PREFIX})
+    file(REMOVE ${FILENAME})
+    find_package(CMakeGet REQUIRED PATHS ${PREFIX})
+endif()
+
+cmake_get_from(${CMAKE_CURRENT_LIST_DIR}/dev-requirements.txt PREFIX ${PREFIX} CMAKE_ARGS -DCMAKE_INSTALL_RPATH=${PREFIX}/lib ${PARSE_UNPARSED_ARGUMENTS})

requirements.txt

-RadeonOpenCompute/rocm-cmake@ec7313aa43de72729c8f66b2e53155e03aa74e20
 google/protobuf
+RadeonOpenCompute/rocm-cmake@3f43e2d493f24abbab4dc189a9ab12cc3ad33baf --build
+ROCmSoftwarePlatform/MIOpen,http://gitlab1.amd.com/pfultz/miopen/-/archive/57dd8372fa043f4bb500ac297f2c479e1c819e89/miopen-57dd8372fa043f4bb500ac297f2c479e1c819e89.tar.bz2 -DMIOPEN_CACHE_DIR=""

src/CMakeLists.txt

@@ -8,3 +8,6 @@ target_include_directories(rtg PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_D
 
 add_subdirectory(onnx)
 add_subdirectory(targets/cpu)
+if(RTG_ENABLE_MIOPEN)
+add_subdirectory(targets/miopen)
+endif()

src/include/rtg/argument.hpp

@@ -26,6 +26,10 @@ 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;
@@ -35,6 +39,12 @@ struct argument : raw_data<argument>
 
     const shape& get_shape() const { return this->m_shape; }
 
+    template <class T>
+    T* cast() const
+    {
+        return reinterpret_cast<T*>(this->data());
+    }
+
     private:
     shape m_shape;
 };

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>
@@ -27,7 +27,7 @@ struct null_deleter
     }
 };
 
-template <class T, class F, F f>
+template <class T, class F, F f> // NOLINT
 using manage_ptr = std::unique_ptr<T, manage_deleter<F, f>>;
 
 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

@@ -235,13 +235,13 @@ struct gemm
     shape compute_shape(std::vector<shape> inputs) const
     {
         check_shapes{inputs}.has(2).same_type();
-        const shape& A = inputs.at(0); 
-        const shape& B = inputs.at(1); 
-        auto t         = A.type();
+        const shape& a = inputs.at(0); 
+        const shape& b = inputs.at(1); 
+        auto t         = a.type();
 
-        if (A.lens()[1] != B.lens()[0])
+        if (a.lens()[1] != b.lens()[0])
             RTG_THROW("Inner dimensions do not match");
-        return {t, {A.lens()[0], B.lens()[1]}};
+        return {t, {a.lens()[0], b.lens()[1]}};
     }
   
     argument compute(shape, std::vector<argument>) const { RTG_THROW("not computable"); }
@@ -446,9 +446,16 @@ struct div
     }
 };
 
-struct reduce
+struct outline
 {
-    std::string name() const { return "reduce"; }
+    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

src/include/rtg/program.hpp

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

src/include/rtg/raw_data.hpp

@@ -6,6 +6,12 @@
 
 namespace rtg {
 
+#define RTG_REQUIRES(...) class = typename std::enable_if<(__VA_ARGS__)>::type
+
+struct raw_data_base
+{
+};
+
 /**
  * @brief Provides a base class for common operations with raw buffer
  *
@@ -15,29 +21,8 @@ namespace rtg {
  *
  */
 template <class Derived>
-struct raw_data
+struct raw_data : raw_data_base
 {
-    friend bool operator==(const Derived& x, const Derived& y)
-    {
-        auto&& xshape = x.get_shape();
-        auto&& yshape = y.get_shape();
-        bool result   = x.empty() && y.empty();
-        if(not result && xshape == yshape)
-        {
-            auto&& xbuffer = x.data();
-            auto&& ybuffer = y.data();
-            // TODO: Dont use tensor view for single values
-            xshape.visit_type([&](auto as) {
-                auto xview = make_view(xshape, as.from(xbuffer));
-                auto yview = make_view(yshape, as.from(ybuffer));
-                result     = xview == yview;
-            });
-        }
-        return result;
-    }
-
-    friend bool operator!=(const Derived& x, const Derived& y) { return !(x == y); }
-
     template <class Stream>
     friend Stream& operator<<(Stream& os, const Derived& d)
     {
@@ -94,8 +79,56 @@ struct raw_data
         this->visit_at([&](auto x) { result = x; }, n);
         return result;
     }
+
+    struct auto_cast
+    {
+        const Derived* self;
+        template <class T>
+        operator T()
+        {
+            return self->template at<T>();
+        }
+        template <class T>
+        operator T*()
+        {
+            // TODO: Check type
+            return reinterpret_cast<T*>(self->data());
+        }
+    };
+
+    auto_cast get() const { return {static_cast<const Derived*>(this)}; }
 };
 
+template <class T,
+          class U,
+          RTG_REQUIRES(std::is_base_of<raw_data_base, T>{} && std::is_base_of<raw_data_base, U>{})>
+bool operator==(const T& x, const U& y)
+{
+    auto&& xshape = x.get_shape();
+    auto&& yshape = y.get_shape();
+    bool result   = x.empty() && y.empty();
+    if(not result && xshape == yshape)
+    {
+        auto&& xbuffer = x.data();
+        auto&& ybuffer = y.data();
+        // TODO: Dont use tensor view for single values
+        xshape.visit_type([&](auto as) {
+            auto xview = make_view(xshape, as.from(xbuffer));
+            auto yview = make_view(yshape, as.from(ybuffer));
+            result     = xview == yview;
+        });
+    }
+    return result;
+}
+
+template <class T,
+          class U,
+          RTG_REQUIRES(std::is_base_of<raw_data_base, T>{} && std::is_base_of<raw_data_base, U>{})>
+bool operator!=(const T& x, const U& y)
+{
+    return !(x == y);
+}
+
 namespace detail {
 template <class V, class... Ts>
 void visit_all_impl(const shape& s, V&& v, Ts&&... xs)

src/include/rtg/shape.hpp

@@ -30,12 +30,15 @@ struct shape
 #define RTG_SHAPE_ENUM_TYPES(x, t) x,
     enum type_t
     {
+        any_type,
         RTG_SHAPE_VISIT_TYPES(RTG_SHAPE_ENUM_TYPES)
     };
 #undef RTG_SHAPE_ENUM_TYPES
 
     template <class T, class = void>
-    struct get_type;
+    struct get_type : std::integral_constant<type_t, any_type>
+    {
+    };
 #define RTG_SHAPE_GET_TYPE(x, t)                              \
     template <class T>                                        \
     struct get_type<t, T> : std::integral_constant<type_t, x> \
@@ -112,6 +115,7 @@ struct shape
     {
         switch(this->m_type)
         {
+        case any_type: RTG_THROW("Cannot visit the any_type");
 #define RTG_SHAPE_VISITOR_CASE(x, t) \
     case x: v(as<t>()); return;
             RTG_SHAPE_VISIT_TYPES(RTG_SHAPE_VISITOR_CASE)

src/include/rtg/tensor_view.hpp

@@ -104,22 +104,6 @@ struct tensor_view
             return m_data + this->size();
     }
 
-    friend bool operator==(const tensor_view<T>& x, const tensor_view<T>& y)
-    {
-        if(x.m_shape == y.m_shape)
-        {
-            for(std::size_t i = 0; i < x.m_shape.elements(); i++)
-            {
-                if(!float_equal(x[i], y[i]))
-                    return false;
-            }
-            return true;
-        }
-        return false;
-    }
-
-    friend bool operator!=(const tensor_view<T>& x, const tensor_view<T>& y) { return !(x == y); }
-
     friend std::ostream& operator<<(std::ostream& os, const tensor_view<T>& x)
     {
         if(!x.empty())
@@ -138,6 +122,27 @@ struct tensor_view
     shape m_shape;
 };
 
+template <class T, class U>
+bool operator==(const tensor_view<T>& x, const tensor_view<U>& y)
+{
+    if(x.get_shape() == y.get_shape())
+    {
+        for(std::size_t i = 0; i < x.get_shape().elements(); i++)
+        {
+            if(!float_equal(x[i], y[i]))
+                return false;
+        }
+        return true;
+    }
+    return false;
+}
+
+template <class T, class U>
+bool operator!=(const tensor_view<T>& x, const tensor_view<U>& y)
+{
+    return !(x == y);
+}
+
 template <class T>
 tensor_view<T> make_view(shape s, T* data)
 {

src/program.cpp

@@ -51,14 +51,39 @@ 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
@@ -87,7 +112,7 @@ void program::compile(const target& t)
         RTG_THROW("Invalid program from compilation");
 }
 
-literal program::eval(std::unordered_map<std::string, argument> params) const
+argument program::eval(std::unordered_map<std::string, argument> params) const
 {
     assert(this->validate() != impl->instructions.end());
     std::unordered_map<const instruction*, argument> results;
@@ -102,6 +127,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());
@@ -113,7 +142,7 @@ literal program::eval(std::unordered_map<std::string, argument> params) const
         }
         results.emplace(std::addressof(ins), result);
     }
-    return literal{result.get_shape(), result.data()};
+    return result;
 }
 
 std::ostream& operator<<(std::ostream& os, const program& p)

src/shape.cpp

@@ -91,6 +91,7 @@ std::string shape::type_string() const
 {
     switch(this->m_type)
     {
+    case any_type: return "any";
 #define RTG_SHAPE_TYPE_STRING_CASE(x, t) \
     case x: return #x;
         RTG_SHAPE_VISIT_TYPES(RTG_SHAPE_TYPE_STRING_CASE)

src/targets/cpu/cpu_target.cpp

@@ -8,7 +8,7 @@ namespace rtg {
 namespace cpu {
 
 template <typename T>
-T zero(const T& x) { return T(0); }
+T zero(const T&) { return T(0); }
 
 struct cpu_convolution
 {
@@ -20,8 +20,6 @@ struct cpu_convolution
     {
         argument result{output_shape};
         visit_all(result, args[0], args[1])([&](auto output, auto input, auto weights) {
-            auto in_n = input.get_shape().lens()[0];
-            auto in_c = input.get_shape().lens()[1];
             auto in_h = input.get_shape().lens()[2];
             auto in_w = input.get_shape().lens()[3];
 
@@ -29,7 +27,10 @@ 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];
@@ -62,25 +63,25 @@ struct cpu_gemm
     argument compute(shape output_shape, std::vector<argument> args) const 
     {
         argument result{output_shape};
-        visit_all(result, args[0], args[1])([&](auto C, auto A, auto B) {
-            auto M = A.get_shape().lens()[0];
-            auto N = B.get_shape().lens()[1];
-            auto K = B.get_shape().lens()[0];
-
-            auto a = A.data();
-            auto b = B.data();
-            auto c = C.data();
-            for (int ii = 0; ii < M; ii++) {
-              for (int jj = 0; jj < N; jj++) {
-                c[ii*N+jj] = 0;
-              }
-            }
-            for (int ii = 0; ii < M; ii++) {
-              for (int kk = 0; kk < K; kk++) {
-                auto aik = a[ii*K+kk];
-                auto* bkj = &b[kk*N];
-                auto* cij = &c[ii*N];
-                for (int jj = 0; jj < N; jj++, cij++, bkj++) {
+        visit_all(result, args[0], args[1])([&](auto cmat, auto amat, auto bmat) {
+            auto m = amat.get_shape().lens()[0];
+            auto n = bmat.get_shape().lens()[1];
+            auto k = bmat.get_shape().lens()[0];
+
+            auto a = amat.data();
+            auto b = bmat.data();
+            auto c = cmat.data();
+            for (int ii = 0; ii < m; ii++) {
+              for (int jj = 0; jj < n; jj++) {
+                c[ii*n+jj] = 0;
+              }
+            }
+            for (int ii = 0; ii < m; ii++) {
+              for (int kk = 0; kk < k; kk++) {
+                auto aik = a[ii*k+kk];
+                auto* bkj = &b[kk*n];
+                auto* cij = &c[ii*n];
+                for (int jj = 0; jj < n; jj++, cij++, bkj++) {
                   *cij += aik*(*bkj);
                 }
               }
@@ -216,10 +217,6 @@ struct softmax2d
                       for (int c = 0; c < nc; c++) {
                           output(b, c, i, j) = output(b, c, i, j)/sum;
                       }
-
-//                       for (int c = 0; c < nc; c++) {
-//                           output(b, c, i, j) = input(b, c, i, j);
-//                       }
                   }
               }
          } 
@@ -348,55 +345,46 @@ struct cpu_apply
 
     void apply_identity(instruction_ref ins)
     {
-        auto&& op = any_cast<identity>(ins->op);
         prog->replace_instruction(ins, cpu_unary<identity_op>{}, ins->arguments);
     }
 
     void apply_softmax(instruction_ref ins)
     {
-        auto&& op = any_cast<softmax>(ins->op);
         prog->replace_instruction(ins, softmax2d{}, ins->arguments);
     }
 
     void apply_tanh(instruction_ref ins)
     {
-        auto&& op = any_cast<tanh>(ins->op);
         prog->replace_instruction(ins, cpu_unary<tanh_op>{}, ins->arguments);
     }
 
     void apply_sigmoid(instruction_ref ins)
     {
-        auto&& op = any_cast<sigmoid>(ins->op);
         prog->replace_instruction(ins, cpu_unary<sigmoid_op>{}, ins->arguments);
     }
 
     void apply_exp(instruction_ref ins)
     {
-        auto&& op = any_cast<exp>(ins->op);
         prog->replace_instruction(ins, cpu_unary<exp_op>{}, ins->arguments);
     }
 
     void apply_neg(instruction_ref ins)
     {
-        auto&& op = any_cast<neg>(ins->op);
         prog->replace_instruction(ins, cpu_unary<neg_op>{}, ins->arguments);
     }
 
     void apply_sin(instruction_ref ins)
     {
-        auto&& op = any_cast<sin>(ins->op);
         prog->replace_instruction(ins, cpu_unary<sin_op>{}, ins->arguments);
     }
 
     void apply_cos(instruction_ref ins)
     {
-        auto&& op = any_cast<cos>(ins->op);
         prog->replace_instruction(ins, cpu_unary<cos_op>{}, ins->arguments);
     }
 
     void apply_tan(instruction_ref ins)
     {
-        auto&& op = any_cast<tan>(ins->op);
         prog->replace_instruction(ins, cpu_unary<tan_op>{}, ins->arguments);
     }
 };

src/targets/miopen/CMakeLists.txt

+
+list(APPEND CMAKE_PREFIX_PATH /opt/rocm /opt/rocm/hip /opt/rocm/hcc)
+find_package(miopen)
+
+if(NOT TARGET MIOpen)
+    message(SEND_ERROR "Cant find miopen")
+endif()
+
+add_library(rtg_miopen
+    miopen_target.cpp
+)
+rocm_clang_tidy_check(rtg_miopen)
+target_link_libraries(rtg_miopen rtg MIOpen)
+target_include_directories(rtg_miopen PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)