merge develop branch.

src/include/migraphx/op/tan.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_TAN_HPP
+#define MIGRAPHX_GUARD_OPERATORS_TAN_HPP
+
+#include <array>
+#include <migraphx/op/unary.hpp>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/streamutils.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/config.hpp>
+#include <cmath>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct tan : unary
+{
+    std::string name() const { return "tan"; }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/op/tanh.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_TANH_HPP
+#define MIGRAPHX_GUARD_OPERATORS_TANH_HPP
+
+#include <array>
+#include <migraphx/op/unary.hpp>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/streamutils.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/config.hpp>
+#include <cmath>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct tanh : unary
+{
+    std::string name() const { return "tanh"; }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/op/transpose.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_TRANSPOSE_HPP
+#define MIGRAPHX_GUARD_OPERATORS_TRANSPOSE_HPP
+
+#include <array>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/streamutils.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/config.hpp>
+#include <cmath>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct transpose
+{
+    std::vector<int64_t> dims;
+
+    template <class Self, class F>
+    static auto reflect(Self& self, F f)
+    {
+        return pack(f(self.dims, "dims"));
+    }
+
+    std::string name() const { return "transpose"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs, *this}.has(1);
+        auto input         = inputs.at(0);
+        auto input_lens    = input.lens();
+        auto input_strides = input.strides();
+        auto t             = input.type();
+        if(dims.size() != input_lens.size())
+        {
+            MIGRAPHX_THROW("Permutation has wrong number of axes");
+        }
+        std::vector<int64_t> axes(dims.size());
+        std::iota(axes.begin(), axes.end(), 0);
+        if(!std::is_permutation(axes.begin(), axes.end(), dims.begin()))
+        {
+            MIGRAPHX_THROW("Invalid permutation");
+        }
+        std::vector<size_t> output_lens(input_lens.size());
+        std::vector<size_t> output_strides(input_lens.size());
+        for(std::size_t i = 0; i < output_lens.size(); i++)
+        {
+            output_lens[i]    = input_lens[dims[i]];
+            output_strides[i] = input_strides[dims[i]];
+        }
+        return {t, output_lens, output_strides};
+    }
+    argument compute(shape output_shape, std::vector<argument> args) const
+    {
+        return {std::move(output_shape), std::move(args.front().data)};
+    }
+    int output_alias(const std::vector<shape>&) const { return 0; }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/op/unary.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_UNARY_HPP
+#define MIGRAPHX_GUARD_OPERATORS_UNARY_HPP
+
+#include <array>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/streamutils.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/config.hpp>
+#include <cmath>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct unary
+{
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return inputs.at(0);
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/op/unsqueeze.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_UNSQUEEZE_HPP
+#define MIGRAPHX_GUARD_OPERATORS_UNSQUEEZE_HPP
+
+#include <array>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/stringutils.hpp>
+#include <migraphx/streamutils.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <migraphx/config.hpp>
+#include <cmath>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct unsqueeze
+{
+    std::vector<int64_t> axes;
+
+    template <class Self, class F>
+    static auto reflect(Self& self, F f)
+    {
+        return pack(f(self.axes, "axes"));
+    }
+
+    std::string name() const { return "unsqueeze"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        auto input_shape     = inputs[0];
+        auto type            = input_shape.type();
+        auto old_lens        = input_shape.lens();
+        std::size_t new_size = old_lens.size() + axes.size();
+        std::vector<std::size_t> new_lens(new_size);
+        std::size_t p = 0;
+        for(std::size_t i = 0; i < new_size; i++)
+        {
+            if(std::find(axes.begin(), axes.end(), i) != axes.end())
+            {
+                new_lens[i] = 1;
+            }
+            else
+            {
+                new_lens[i] = old_lens[p++];
+            }
+        }
+        return shape{type, new_lens};
+    }
+    argument compute(shape output_shape, std::vector<argument> args) const
+    {
+        return {std::move(output_shape), std::move(args.front().data)};
+    }
+    int output_alias(const std::vector<shape>&) const { return 0; }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/rewrite_rnn.hpp

@@ -4,7 +4,7 @@
 #include <string>
 #include <vector>
 #include <migraphx/instruction_ref.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/operation.hpp>
 #include <migraphx/config.hpp>
 
 namespace migraphx {

src/opt/memory_coloring_impl.cpp

+#include <migraphx/op/load.hpp>
 #include "memory_coloring_impl.hpp"
 
 namespace migraphx {

src/opt/memory_coloring_impl.hpp

@@ -3,7 +3,6 @@
 #include <migraphx/program.hpp>
 #include <migraphx/stringutils.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/pass_config.hpp>
 #include <migraphx/config.hpp>

src/program.cpp

 #include <migraphx/program.hpp>
 #include <migraphx/stringutils.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/identity.hpp>
 #include <migraphx/target.hpp>
 #include <migraphx/env.hpp>
 #include <migraphx/ranges.hpp>

src/schedule.cpp

 #include <migraphx/schedule.hpp>
 #include <migraphx/program.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/identity.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/dfor.hpp>
 #include <migraphx/functional.hpp>

src/simplify_algebra.cpp

 #include <migraphx/simplify_algebra.hpp>
 #include <migraphx/program.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/add.hpp>
 #include <migraphx/matcher.hpp>
 #include <migraphx/literal.hpp>
 

src/simplify_reshapes.cpp

 #include <migraphx/simplify_reshapes.hpp>
 #include <migraphx/program.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/as_shape.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/ranges.hpp>
 #include <unordered_set>

src/targets/gpu/eliminate_workspace.cpp

@@ -2,7 +2,6 @@
 #include <migraphx/gpu/hip.hpp>
 #include <migraphx/program.hpp>
 #include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/ranges.hpp>
 #include <migraphx/stringutils.hpp>

src/targets/gpu/gemm.cpp

@@ -179,6 +179,7 @@ argument miopen_gemm::compute(context& ctx,
                               const std::vector<argument>& args) const
 {
     bool is_3inputs = (args.size() == 4);
+    float beta = 0.0f;
     if(is_3inputs)
     {
         output_shape.visit_type([&](auto as) {
@@ -189,61 +190,15 @@ argument miopen_gemm::compute(context& ctx,
                            hipMemcpyDeviceToDevice,
                            ctx.get_stream().get());
         });
-
-        output_shape.visit_type([&](auto as) {
-            auto n_dim        = output_shape.lens().size();
-            auto dim_1        = n_dim - 1;
-            auto dim_0        = n_dim - 2;
-            auto alpha_r      = to_rocblas_type(as(op.alpha));
-            auto beta_r       = to_rocblas_type(as(op.beta));
-            bool transa       = args[0].get_shape().transposed();
-            bool transb       = args[1].get_shape().transposed();
-            rocblas_int lda   = args[0].get_shape().strides()[transa ? dim_1 : dim_0];
-            rocblas_int ldb   = args[1].get_shape().strides()[transb ? dim_1 : dim_0];
-            rocblas_int ldc   = args[3].get_shape().strides()[dim_0];
-            auto out_lens     = output_shape.lens();
-            rocblas_int m     = out_lens[dim_0];
-            rocblas_int n     = out_lens[dim_1];
-            rocblas_int k     = args[0].get_shape().lens()[dim_1];
-            auto num_matrices = std::accumulate(out_lens.rbegin() + 2,
-                                                out_lens.rend(),
-                                                std::size_t{1},
-                                                std::multiplies<std::size_t>());
-            auto to_pointer   = [&](auto&& arg) { return to_rocblas_type(as.from(arg.data())); };
-            generic_rocblas_batched_gemm(
-                as,
-                ctx.get_stream().get_rocblas(),
-                transb ? rocblas_operation_transpose : rocblas_operation_none,
-                transa ? rocblas_operation_transpose : rocblas_operation_none,
-                n,
-                m,
-                k,
-                &alpha_r,
-                to_pointer(args[1]),
-                ldb,
-                k * n,
-                to_pointer(args[0]),
-                lda,
-                m * k,
-                &beta_r,
-                to_pointer(args[3]),
-                ldc,
-                m * n,
-                num_matrices);
-        });
-
-        return args[3];
+        beta = op.beta;
     }
 
-    // 2 input argument cases
-    // vector inner product
     auto a_lens = args[0].get_shape().lens();
     auto b_lens = args[1].get_shape().lens();
     output_shape.visit_type([&](auto as) {
         auto n_dim        = output_shape.lens().size();
         auto dim_1        = n_dim - 1;
         auto dim_0        = n_dim - 2;
-        float beta        = 0.0f;
         auto alpha_r      = to_rocblas_type(as(op.alpha));
         auto beta_r       = to_rocblas_type(as(beta));
         bool transa       = args[0].get_shape().transposed();
@@ -273,7 +228,7 @@ argument miopen_gemm::compute(context& ctx,
                                  to_pointer(args[0]),
                                  lda,
                                  &beta_r,
-                                 to_pointer(args[2]),
+                                 (is_3inputs ? to_pointer(args[3]) : to_pointer(args[2])),
                                  ldc);
         }
         else
@@ -294,14 +249,14 @@ argument miopen_gemm::compute(context& ctx,
                 lda,
                 m * k,
                 &beta_r,
-                to_pointer(args[2]),
+                (is_3inputs ? to_pointer(args[3]) : to_pointer(args[2])),
                 ldc,
                 m * n,
                 num_matrices);
         }
     });
 
-    return args[2];
+    return (is_3inputs ? args[3] : args[2]);
 }
 
 } // namespace gpu

src/targets/gpu/include/migraphx/gpu/batchnorm.hpp

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_BATCHNORM_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/batch_norm.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/include/migraphx/gpu/concat.hpp

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_CONCAT_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/concat.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/include/migraphx/gpu/contiguous.hpp

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_CONTIGUOUS_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/contiguous.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/include/migraphx/gpu/convolution.hpp

@@ -2,7 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_CONVOLUTION_HPP
 
 #include <migraphx/shape.hpp>
-#include <migraphx/operators.hpp>
+#include <migraphx/op/convolution.hpp>
 #include <migraphx/gpu/miopen.hpp>
 
 namespace migraphx {

src/targets/gpu/include/migraphx/gpu/gather.hpp

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_GATHER_HPP
 
 #include <migraphx/shape.hpp>
+#include <migraphx/op/gather.hpp>
 #include <migraphx/gpu/miopen.hpp>
 
 namespace migraphx {